From: Vladislav Vinogradov <no@email>
Date: Wed, 17 Aug 2011 13:46:36 +0000 (+0000)
Subject: fixed bug in gpu::cvtColor
X-Git-Tag: accepted/2.0/20130307.220821~1991
X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=971e35f2837bff287ae0732625c381e1760f7312;p=profile%2Fivi%2Fopencv.git

fixed bug in gpu::cvtColor
---

diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp
index c539f3a..0f42d4d 100644
--- a/modules/gpu/src/color.cpp
+++ b/modules/gpu/src/color.cpp
@@ -51,7 +51,7 @@ void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, Stream&) { throw_nogpu(
 
 #else /* !defined (HAVE_CUDA) */
 
-namespace cv { namespace gpu {  namespace color  
+namespace cv { namespace gpu {  namespace device  
 {
     #define OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name) \
         void name(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream);
@@ -207,7 +207,7 @@ namespace
 
     void bgr_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {bgr_to_rgb_8u, 0, bgr_to_rgb_16u, 0, 0, bgr_to_rgb_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -220,7 +220,7 @@ namespace
 
     void bgr_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {bgr_to_bgra_8u, 0, bgr_to_bgra_16u, 0, 0, bgr_to_bgra_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -233,7 +233,7 @@ namespace
 
     void bgr_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {bgr_to_rgba_8u, 0, bgr_to_rgba_16u, 0, 0, bgr_to_rgba_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -246,7 +246,7 @@ namespace
 
     void bgra_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {bgra_to_bgr_8u, 0, bgra_to_bgr_16u, 0, 0, bgra_to_bgr_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -259,7 +259,7 @@ namespace
 
     void bgra_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {bgra_to_rgb_8u, 0, bgra_to_rgb_16u, 0, 0, bgra_to_rgb_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -272,7 +272,7 @@ namespace
 
     void bgra_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {bgra_to_rgba_8u, 0, bgra_to_rgba_16u, 0, 0, bgra_to_rgba_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -290,7 +290,7 @@ namespace
 
         dst.create(src.size(), CV_8UC2);        
 
-        color::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream));
+        device::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -300,7 +300,7 @@ namespace
 
         dst.create(src.size(), CV_8UC2);        
 
-        color::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream));
+        device::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream));
     }
 
     void rgb_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -310,7 +310,7 @@ namespace
 
         dst.create(src.size(), CV_8UC2);        
 
-        color::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream));
+        device::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream));
     }
 
     void rgb_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -320,7 +320,7 @@ namespace
 
         dst.create(src.size(), CV_8UC2);        
 
-        color::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream));
+        device::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgra_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -330,7 +330,7 @@ namespace
 
         dst.create(src.size(), CV_8UC2);        
 
-        color::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream));
+        device::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgra_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -340,7 +340,7 @@ namespace
 
         dst.create(src.size(), CV_8UC2);        
 
-        color::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream));
+        device::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream));
     }
 
     void rgba_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -350,7 +350,7 @@ namespace
 
         dst.create(src.size(), CV_8UC2);        
 
-        color::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream));
+        device::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream));
     }
 
     void rgba_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -360,7 +360,7 @@ namespace
 
         dst.create(src.size(), CV_8UC2);        
 
-        color::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream));
+        device::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr555_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -370,7 +370,7 @@ namespace
 
         dst.create(src.size(), CV_8UC3);        
 
-        color::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream));
+        device::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr565_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -380,7 +380,7 @@ namespace
 
         dst.create(src.size(), CV_8UC3);        
 
-        color::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream));
+        device::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr555_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -390,7 +390,7 @@ namespace
 
         dst.create(src.size(), CV_8UC3);        
 
-        color::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream));
+        device::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr565_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -400,7 +400,7 @@ namespace
 
         dst.create(src.size(), CV_8UC3);        
 
-        color::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream));
+        device::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr555_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -410,7 +410,7 @@ namespace
 
         dst.create(src.size(), CV_8UC4);        
 
-        color::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream));
+        device::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr565_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -420,7 +420,7 @@ namespace
 
         dst.create(src.size(), CV_8UC4);        
 
-        color::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream));
+        device::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr555_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -430,7 +430,7 @@ namespace
 
         dst.create(src.size(), CV_8UC4);        
 
-        color::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream));
+        device::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr565_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -440,12 +440,12 @@ namespace
 
         dst.create(src.size(), CV_8UC4);        
 
-        color::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream));
+        device::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream));
     }
 
     void gray_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {gray_to_bgr_8u, 0, gray_to_bgr_16u, 0, 0, gray_to_bgr_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -458,7 +458,7 @@ namespace
 
     void gray_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {gray_to_bgra_8u, 0, gray_to_bgra_16u, 0, 0, gray_to_bgra_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -476,7 +476,7 @@ namespace
 
         dst.create(src.size(), CV_8UC2);        
 
-        color::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream));
+        device::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream));
     }
 
     void gray_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -486,7 +486,7 @@ namespace
 
         dst.create(src.size(), CV_8UC2);        
 
-        color::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream));
+        device::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr555_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -496,7 +496,7 @@ namespace
 
         dst.create(src.size(), CV_8UC1);        
 
-        color::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream));
+        device::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream));
     }
 
     void bgr565_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)
@@ -506,12 +506,12 @@ namespace
 
         dst.create(src.size(), CV_8UC1);        
 
-        color::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream));
+        device::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream));
     }
 
     void rgb_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {rgb_to_gray_8u, 0, rgb_to_gray_16u, 0, 0, rgb_to_gray_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -524,7 +524,7 @@ namespace
 
     void bgr_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {bgr_to_gray_8u, 0, bgr_to_gray_16u, 0, 0, bgr_to_gray_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -537,7 +537,7 @@ namespace
 
     void rgba_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {rgba_to_gray_8u, 0, rgba_to_gray_16u, 0, 0, rgba_to_gray_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -550,7 +550,7 @@ namespace
 
     void bgra_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[] = {bgra_to_gray_8u, 0, bgra_to_gray_16u, 0, 0, bgra_to_gray_32f};
         
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);
@@ -563,7 +563,7 @@ namespace
     
     void rgb_to_yuv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -589,7 +589,7 @@ namespace
 
     void bgr_to_yuv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -615,7 +615,7 @@ namespace
 
     void yuv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -641,7 +641,7 @@ namespace
 
     void yuv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -667,7 +667,7 @@ namespace
     
     void rgb_to_YCrCb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -693,7 +693,7 @@ namespace
 
     void bgr_to_YCrCb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -719,7 +719,7 @@ namespace
 
     void YCrCb_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -745,7 +745,7 @@ namespace
 
     void YCrCb_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -771,7 +771,7 @@ namespace
 
     void rgb_to_xyz(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -797,7 +797,7 @@ namespace
 
     void bgr_to_xyz(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -823,7 +823,7 @@ namespace
 
     void xyz_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -849,7 +849,7 @@ namespace
 
     void xyz_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -875,7 +875,7 @@ namespace
 
     void rgb_to_hsv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -901,7 +901,7 @@ namespace
 
     void bgr_to_hsv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -927,7 +927,7 @@ namespace
 
     void hsv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -953,7 +953,7 @@ namespace
 
     void hsv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -979,7 +979,7 @@ namespace
 
     void rgb_to_hls(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1005,7 +1005,7 @@ namespace
 
     void bgr_to_hls(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1031,7 +1031,7 @@ namespace
 
     void hls_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1057,7 +1057,7 @@ namespace
 
     void hls_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1083,7 +1083,7 @@ namespace
 
     void rgb_to_hsv_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1109,7 +1109,7 @@ namespace
 
     void bgr_to_hsv_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1135,7 +1135,7 @@ namespace
 
     void hsv_to_rgb_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1161,7 +1161,7 @@ namespace
 
     void hsv_to_bgr_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1187,7 +1187,7 @@ namespace
 
     void rgb_to_hls_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1213,7 +1213,7 @@ namespace
 
     void bgr_to_hls_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1239,7 +1239,7 @@ namespace
 
     void hls_to_rgb_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1265,7 +1265,7 @@ namespace
 
     void hls_to_bgr_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
     {
-        using namespace cv::gpu::color;
+        using namespace cv::gpu::device;
         static const gpu_func_t funcs[2][2][6] = 
         {
             {
@@ -1364,7 +1364,7 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream
         rgb_to_hls,             // CV_RGB2HLS     =53
 
         hsv_to_bgr,             // CV_HSV2BGR     =54
-        bgr_to_rgb,             // CV_HSV2RGB     =55
+        hsv_to_rgb,             // CV_HSV2RGB     =55
 
         0,                      // CV_Lab2BGR     =56
         0,                      // CV_Lab2RGB     =57
diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu
index 2755fff..a16c82c 100644
--- a/modules/gpu/src/cuda/color.cu
+++ b/modules/gpu/src/cuda/color.cu
@@ -44,11 +44,15 @@
 #include "opencv2/gpu/device/transform.hpp"
 #include "opencv2/gpu/device/color.hpp"
 
-using namespace cv::gpu;
-using namespace cv::gpu::device;
-
-namespace cv { namespace gpu { namespace color
+namespace cv { namespace gpu { namespace device
 {
+    template <> struct TransformFunctorTraits<bgra_to_rgba_traits<uchar>::functor_type> : DefaultTransformFunctorTraits<bgra_to_rgba_traits<uchar>::functor_type>
+    {
+        enum { smart_block_dim_x = 8 };
+        enum { smart_block_dim_y = 8 };
+        enum { smart_shift = 4 };
+    };
+
     #define OPENCV_GPU_IMPLEMENT_CVTCOLOR(name, traits) \
         void name(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream) \
         { \
diff --git a/modules/gpu/src/opencv2/gpu/device/detail/color.hpp b/modules/gpu/src/opencv2/gpu/device/detail/color.hpp
index 748769e..467a0ca 100644
--- a/modules/gpu/src/opencv2/gpu/device/detail/color.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/detail/color.hpp
@@ -100,15 +100,32 @@ namespace cv { namespace gpu { namespace device
 
     namespace detail
     {
-        template <typename T, typename D, int bidx> struct RGB2RGB : public unary_function<T, D>
+        template <typename T, int scn, int dcn, int bidx> struct RGB2RGB : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
         {
-            __device__ D operator()(const T& src) const
+            __device__ typename TypeVec<T, dcn>::vec_type operator()(const typename TypeVec<T, scn>::vec_type& src) const
             {
-                D dst;
+                typename TypeVec<T, dcn>::vec_type dst;
+
                 dst.x = (&src.x)[bidx];
                 dst.y = src.y;
                 dst.z = (&src.x)[bidx^2];
-                setAlpha(dst, getAlpha<typename VecTraits<T>::elem_type>(src));
+                setAlpha(dst, getAlpha<T>(src));
+
+                return dst;
+            }
+        };
+
+        template <> struct RGB2RGB<uchar, 4, 4, 2> : unary_function<uint, uint>
+        {
+            __device__ uint operator()(uint src) const
+            {
+                uint dst = 0;
+
+                dst |= (0xff & (src >> 16));
+                dst |= (0xff & (src >> 8)) << 8;
+                dst |= (0xff & (src)) << 16;
+                dst |= (0xff & (src >> 24)) << 24;
+
                 return dst;
             }
         };
@@ -117,10 +134,10 @@ namespace cv { namespace gpu { namespace device
 #define OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(name, scn, dcn, bidx) \
     template <typename T> struct name ## _traits \
     { \
-        typedef detail::RGB2RGB<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type, bidx> functor_type; \
+        typedef detail::RGB2RGB<T, scn, dcn, bidx> functor_type; \
         static __host__ __device__ __forceinline__ functor_type create_functor() \
         { \
-            return detail::RGB2RGB<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type, bidx>(); \
+            return functor_type(); \
         } \
     };
 
diff --git a/modules/gpu/test/test_imgproc.cpp b/modules/gpu/test/test_imgproc.cpp
index cc39a13..9f33783 100644
--- a/modules/gpu/test/test_imgproc.cpp
+++ b/modules/gpu/test/test_imgproc.cpp
@@ -561,44 +561,904 @@ TEST_P(CvtColor, BGR2RGBA)
     EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
 }
 
+TEST_P(CvtColor, BGR2BGRA)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src = img;
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR2BGRA);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2BGRA);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
 TEST_P(CvtColor, BGRA2RGB)
 {
     PRINT_PARAM(devInfo);
     PRINT_TYPE(type);
 
     cv::Mat src;
-    cv::cvtColor(img, src, CV_BGR2BGRA);
+    cv::cvtColor(img, src, CV_BGR2BGRA);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGRA2RGB);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2RGB);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGRA2BGR)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGRA);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGRA2BGR);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2BGR);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGRA2RGBA)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGRA);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGRA2RGBA);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2RGBA);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR2GRAY)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src = img;
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR2GRAY);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2GRAY);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, RGB2GRAY)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2RGB);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_RGB2GRAY);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2GRAY);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, GRAY2BGR)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2GRAY);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_GRAY2BGR);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_GRAY2BGR);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, GRAY2BGRA)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2GRAY);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_GRAY2BGRA, 4);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_GRAY2BGRA, 4);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGRA2GRAY)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGRA);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGRA2GRAY);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2GRAY);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, RGBA2GRAY)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2RGBA);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_RGBA2GRAY);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGBA2GRAY);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, BGR2BGR565)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src = img;
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR2BGR565);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2BGR565);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, RGB2BGR565)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2RGB);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_RGB2BGR565);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2BGR565);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR5652BGR)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGR565);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR5652BGR);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5652BGR);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR5652RGB)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGR565);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR5652RGB);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5652RGB);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGRA2BGR565)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGRA);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGRA2BGR565);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2BGR565);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, RGBA2BGR565)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2RGBA);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_RGBA2BGR565);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGBA2BGR565);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR5652BGRA)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGR565);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR5652BGRA, 4);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5652BGRA, 4);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR5652RGBA)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGR565);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR5652RGBA, 4);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5652RGBA, 4);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, GRAY2BGR565)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2GRAY);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_GRAY2BGR565);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_GRAY2BGR565);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR5652GRAY)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGR565);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR5652GRAY);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5652GRAY);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR2BGR555)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src = img;
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR2BGR555);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2BGR555);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, RGB2BGR555)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2RGB);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_RGB2BGR555);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2BGR555);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR5552BGR)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGR555);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR5552BGR);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5552BGR);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR5552RGB)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGR555);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR5552RGB);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5552RGB);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGRA2BGR555)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGRA);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGRA2BGR555);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2BGR555);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, RGBA2BGR555)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2RGBA);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_RGBA2BGR555);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGBA2BGR555);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR5552BGRA)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGR555);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR5552BGRA, 4);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5552BGRA, 4);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR5552RGBA)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGR555);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR5552RGBA, 4);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5552RGBA, 4);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, GRAY2BGR555)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2GRAY);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_GRAY2BGR555);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_GRAY2BGR555);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR5552GRAY)
+{
+    if (type != CV_8U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2BGR555);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR5552GRAY);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5552GRAY);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+}
+
+TEST_P(CvtColor, BGR2XYZ)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src = img;
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR2XYZ);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2XYZ);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, RGB2XYZ)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2RGB);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR2XYZ);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2XYZ);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, XYZ2BGR)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2XYZ);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_XYZ2BGR);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_XYZ2BGR);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, XYZ2RGB)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2XYZ);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_XYZ2RGB);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_XYZ2RGB);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, BGR2YCrCb)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src = img;
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_BGR2YCrCb);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2YCrCb);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, RGB2YCrCb)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2RGB);
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_BGRA2RGB);
+    cv::cvtColor(src, dst_gold, CV_RGB2YCrCb);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2RGB);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2YCrCb);
 
         gpuRes.download(dst);
     );
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
 }
 
-TEST_P(CvtColor, BGR2YCrCb)
+TEST_P(CvtColor, YCrCb2BGR)
 {
     PRINT_PARAM(devInfo);
     PRINT_TYPE(type);
 
-    cv::Mat src = img;
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2YCrCb);
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_BGR2YCrCb);
+    cv::cvtColor(src, dst_gold, CV_YCrCb2BGR);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2YCrCb);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_YCrCb2BGR);
 
         gpuRes.download(dst);
     );
@@ -629,97 +1489,109 @@ TEST_P(CvtColor, YCrCb2RGB)
     EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
 }
 
-TEST_P(CvtColor, BGR2YUV)
+TEST_P(CvtColor, BGR2HSV)
 {
+    if (type == CV_16U)
+        return;
+
     PRINT_PARAM(devInfo);
     PRINT_TYPE(type);
 
     cv::Mat src = img;
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_BGR2YUV);
+    cv::cvtColor(src, dst_gold, CV_BGR2HSV);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2YUV);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HSV);
 
         gpuRes.download(dst);
     );
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
-TEST_P(CvtColor, YUV2BGR)
+TEST_P(CvtColor, RGB2HSV)
 {
+    if (type == CV_16U)
+        return;
+
     PRINT_PARAM(devInfo);
     PRINT_TYPE(type);
 
     cv::Mat src;
-    cv::cvtColor(img, src, CV_BGR2YUV);
+    cv::cvtColor(img, src, CV_BGR2RGB);
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_YUV2BGR);
+    cv::cvtColor(src, dst_gold, CV_RGB2HSV);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_YUV2BGR);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2HSV);
 
         gpuRes.download(dst);
     );
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
-TEST_P(CvtColor, BGR2XYZ)
+TEST_P(CvtColor, BGR2HLS)
 {
+    if (type == CV_16U)
+        return;
+
     PRINT_PARAM(devInfo);
     PRINT_TYPE(type);
 
     cv::Mat src = img;
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_BGR2XYZ);
+    cv::cvtColor(src, dst_gold, CV_BGR2HLS);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2XYZ);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HLS);
 
         gpuRes.download(dst);
     );
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
-TEST_P(CvtColor, XYZ2BGR)
+TEST_P(CvtColor, RGB2HLS)
 {
+    if (type == CV_16U)
+        return;
+
     PRINT_PARAM(devInfo);
     PRINT_TYPE(type);
 
     cv::Mat src;
-    cv::cvtColor(img, src, CV_BGR2XYZ);
+    cv::cvtColor(img, src, CV_BGR2RGB);
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_XYZ2BGR);
+    cv::cvtColor(src, dst_gold, CV_RGB2HLS);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_XYZ2BGR);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2HLS);
 
         gpuRes.download(dst);
     );
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
-TEST_P(CvtColor, BGR2HSV)
+TEST_P(CvtColor, HSV2BGR)
 {
     if (type == CV_16U)
         return;
@@ -727,16 +1599,17 @@ TEST_P(CvtColor, BGR2HSV)
     PRINT_PARAM(devInfo);
     PRINT_TYPE(type);
 
-    cv::Mat src = img;
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2HSV);
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_BGR2HSV);
+    cv::cvtColor(src, dst_gold, CV_HSV2BGR);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HSV);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2BGR);
 
         gpuRes.download(dst);
     );
@@ -744,7 +1617,7 @@ TEST_P(CvtColor, BGR2HSV)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
-TEST_P(CvtColor, HSV2BGR)
+TEST_P(CvtColor, HSV2RGB)
 {
     if (type == CV_16U)
         return;
@@ -755,14 +1628,14 @@ TEST_P(CvtColor, HSV2BGR)
     cv::Mat src;
     cv::cvtColor(img, src, CV_BGR2HSV);
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_HSV2BGR);
+    cv::cvtColor(src, dst_gold, CV_HSV2RGB);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2BGR);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2RGB);
 
         gpuRes.download(dst);
     );
@@ -770,7 +1643,7 @@ TEST_P(CvtColor, HSV2BGR)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
-TEST_P(CvtColor, BGR2HSV_FULL)
+TEST_P(CvtColor, HLS2BGR)
 {
     if (type == CV_16U)
         return;
@@ -778,16 +1651,17 @@ TEST_P(CvtColor, BGR2HSV_FULL)
     PRINT_PARAM(devInfo);
     PRINT_TYPE(type);
 
-    cv::Mat src = img;
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2HLS);
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_BGR2HSV_FULL);
+    cv::cvtColor(src, dst_gold, CV_HLS2BGR);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HSV_FULL);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HLS2BGR);
 
         gpuRes.download(dst);
     );
@@ -795,7 +1669,7 @@ TEST_P(CvtColor, BGR2HSV_FULL)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
-TEST_P(CvtColor, HSV2BGR_FULL)
+TEST_P(CvtColor, HLS2RGB)
 {
     if (type == CV_16U)
         return;
@@ -804,16 +1678,16 @@ TEST_P(CvtColor, HSV2BGR_FULL)
     PRINT_TYPE(type);
 
     cv::Mat src;
-    cv::cvtColor(img, src, CV_BGR2HSV_FULL);
+    cv::cvtColor(img, src, CV_BGR2HLS);
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_HSV2BGR_FULL);
+    cv::cvtColor(src, dst_gold, CV_HLS2RGB);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2BGR_FULL);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HLS2RGB);
 
         gpuRes.download(dst);
     );
@@ -821,7 +1695,7 @@ TEST_P(CvtColor, HSV2BGR_FULL)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
-TEST_P(CvtColor, BGR2HLS)
+TEST_P(CvtColor, BGR2HSV_FULL)
 {
     if (type == CV_16U)
         return;
@@ -831,14 +1705,14 @@ TEST_P(CvtColor, BGR2HLS)
 
     cv::Mat src = img;
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_BGR2HLS);
+    cv::cvtColor(src, dst_gold, CV_BGR2HSV_FULL);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HLS);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HSV_FULL);
 
         gpuRes.download(dst);
     );
@@ -846,7 +1720,7 @@ TEST_P(CvtColor, BGR2HLS)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
-TEST_P(CvtColor, HLS2BGR)
+TEST_P(CvtColor, RGB2HSV_FULL)
 {
     if (type == CV_16U)
         return;
@@ -855,16 +1729,16 @@ TEST_P(CvtColor, HLS2BGR)
     PRINT_TYPE(type);
 
     cv::Mat src;
-    cv::cvtColor(img, src, CV_BGR2HLS);
+    cv::cvtColor(img, src, CV_BGR2RGB);
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_HLS2BGR);
+    cv::cvtColor(src, dst_gold, CV_RGB2HSV_FULL);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HLS2BGR);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2HSV_FULL);
 
         gpuRes.download(dst);
     );
@@ -897,6 +1771,84 @@ TEST_P(CvtColor, BGR2HLS_FULL)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
+TEST_P(CvtColor, RGB2HLS_FULL)
+{
+    if (type == CV_16U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2RGB);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_RGB2HLS_FULL);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2HLS_FULL);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
+}
+
+TEST_P(CvtColor, HSV2BGR_FULL)
+{
+    if (type == CV_16U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2HSV_FULL);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_HSV2BGR_FULL);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2BGR_FULL);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
+}
+
+TEST_P(CvtColor, HSV2RGB_FULL)
+{
+    if (type == CV_16U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2HSV_FULL);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_HSV2RGB_FULL);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2RGB_FULL);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
+}
+
 TEST_P(CvtColor, HLS2BGR_FULL)
 {
     if (type == CV_16U)
@@ -923,21 +1875,47 @@ TEST_P(CvtColor, HLS2BGR_FULL)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
 }
 
-TEST_P(CvtColor, BGR2GRAY)
+TEST_P(CvtColor, HLS2RGB_FULL)
+{
+    if (type == CV_16U)
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2HLS_FULL);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_HLS2RGB_FULL);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HLS2RGB_FULL);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);
+}
+
+TEST_P(CvtColor, BGR2YUV)
 {
     PRINT_PARAM(devInfo);
     PRINT_TYPE(type);
 
     cv::Mat src = img;
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_BGR2GRAY);
+    cv::cvtColor(src, dst_gold, CV_BGR2YUV);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2GRAY);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2YUV);
 
         gpuRes.download(dst);
     );
@@ -945,27 +1923,73 @@ TEST_P(CvtColor, BGR2GRAY)
     EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
 }
 
-TEST_P(CvtColor, GRAY2RGB)
+TEST_P(CvtColor, RGB2YUV)
 {
     PRINT_PARAM(devInfo);
     PRINT_TYPE(type);
 
     cv::Mat src;
-    cv::cvtColor(img, src, CV_BGR2GRAY);
+    cv::cvtColor(img, src, CV_BGR2RGB);
     cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, CV_GRAY2RGB);
+    cv::cvtColor(src, dst_gold, CV_RGB2YUV);
 
     cv::Mat dst;
 
     ASSERT_NO_THROW(
         cv::gpu::GpuMat gpuRes;
 
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_GRAY2RGB);
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2YUV);
 
         gpuRes.download(dst);
     );
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, YUV2BGR)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_BGR2YUV);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_YUV2BGR);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_YUV2BGR);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+TEST_P(CvtColor, YUV2RGB)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type);
+
+    cv::Mat src;
+    cv::cvtColor(img, src, CV_RGB2YUV);
+    cv::Mat dst_gold;
+    cv::cvtColor(src, dst_gold, CV_YUV2RGB);
+
+    cv::Mat dst;
+
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuRes;
+
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_YUV2RGB);
+
+        gpuRes.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
 }
 
 INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor, testing::Combine(