fixed bug in gpu::cvtColor

diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp
index c539f3a..0f42d4d 100644 (file)
--- 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(
 \r
 #else /* !defined (HAVE_CUDA) */\r
 \r
-namespace cv { namespace gpu {  namespace color  \r
+namespace cv { namespace gpu {  namespace device  \r
 {\r
     #define OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name) \\r
         void name(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream);\r
@@ -207,7 +207,7 @@ namespace
 \r
     void bgr_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {bgr_to_rgb_8u, 0, bgr_to_rgb_16u, 0, 0, bgr_to_rgb_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -220,7 +220,7 @@ namespace
 \r
     void bgr_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {bgr_to_bgra_8u, 0, bgr_to_bgra_16u, 0, 0, bgr_to_bgra_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -233,7 +233,7 @@ namespace
 \r
     void bgr_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {bgr_to_rgba_8u, 0, bgr_to_rgba_16u, 0, 0, bgr_to_rgba_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -246,7 +246,7 @@ namespace
 \r
     void bgra_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {bgra_to_bgr_8u, 0, bgra_to_bgr_16u, 0, 0, bgra_to_bgr_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -259,7 +259,7 @@ namespace
 \r
     void bgra_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {bgra_to_rgb_8u, 0, bgra_to_rgb_16u, 0, 0, bgra_to_rgb_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -272,7 +272,7 @@ namespace
 \r
     void bgra_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {bgra_to_rgba_8u, 0, bgra_to_rgba_16u, 0, 0, bgra_to_rgba_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -290,7 +290,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC2);        \r
 \r
-        color::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -300,7 +300,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC2);        \r
 \r
-        color::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void rgb_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -310,7 +310,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC2);        \r
 \r
-        color::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream));\r
+        device::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void rgb_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -320,7 +320,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC2);        \r
 \r
-        color::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream));\r
+        device::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgra_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -330,7 +330,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC2);        \r
 \r
-        color::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgra_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -340,7 +340,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC2);        \r
 \r
-        color::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void rgba_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -350,7 +350,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC2);        \r
 \r
-        color::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream));\r
+        device::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void rgba_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -360,7 +360,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC2);        \r
 \r
-        color::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream));\r
+        device::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr555_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -370,7 +370,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC3);        \r
 \r
-        color::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr565_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -380,7 +380,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC3);        \r
 \r
-        color::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr555_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -390,7 +390,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC3);        \r
 \r
-        color::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr565_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -400,7 +400,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC3);        \r
 \r
-        color::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr555_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -410,7 +410,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC4);        \r
 \r
-        color::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr565_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -420,7 +420,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC4);        \r
 \r
-        color::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr555_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -430,7 +430,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC4);        \r
 \r
-        color::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr565_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -440,12 +440,12 @@ namespace
 \r
         dst.create(src.size(), CV_8UC4);        \r
 \r
-        color::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void gray_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {gray_to_bgr_8u, 0, gray_to_bgr_16u, 0, 0, gray_to_bgr_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -458,7 +458,7 @@ namespace
 \r
     void gray_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {gray_to_bgra_8u, 0, gray_to_bgra_16u, 0, 0, gray_to_bgra_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -476,7 +476,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC2);        \r
 \r
-        color::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream));\r
+        device::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void gray_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -486,7 +486,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC2);        \r
 \r
-        color::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream));\r
+        device::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr555_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -496,7 +496,7 @@ namespace
 \r
         dst.create(src.size(), CV_8UC1);        \r
 \r
-        color::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void bgr565_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
@@ -506,12 +506,12 @@ namespace
 \r
         dst.create(src.size(), CV_8UC1);        \r
 \r
-        color::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream));\r
+        device::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream));\r
     }\r
 \r
     void rgb_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {rgb_to_gray_8u, 0, rgb_to_gray_16u, 0, 0, rgb_to_gray_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -524,7 +524,7 @@ namespace
 \r
     void bgr_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {bgr_to_gray_8u, 0, bgr_to_gray_16u, 0, 0, bgr_to_gray_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -537,7 +537,7 @@ namespace
 \r
     void rgba_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {rgba_to_gray_8u, 0, rgba_to_gray_16u, 0, 0, rgba_to_gray_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -550,7 +550,7 @@ namespace
 \r
     void bgra_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[] = {bgra_to_gray_8u, 0, bgra_to_gray_16u, 0, 0, bgra_to_gray_32f};\r
         \r
         CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F);\r
@@ -563,7 +563,7 @@ namespace
     \r
     void rgb_to_yuv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -589,7 +589,7 @@ namespace
 \r
     void bgr_to_yuv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -615,7 +615,7 @@ namespace
 \r
     void yuv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -641,7 +641,7 @@ namespace
 \r
     void yuv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -667,7 +667,7 @@ namespace
     \r
     void rgb_to_YCrCb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -693,7 +693,7 @@ namespace
 \r
     void bgr_to_YCrCb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -719,7 +719,7 @@ namespace
 \r
     void YCrCb_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -745,7 +745,7 @@ namespace
 \r
     void YCrCb_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -771,7 +771,7 @@ namespace
 \r
     void rgb_to_xyz(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -797,7 +797,7 @@ namespace
 \r
     void bgr_to_xyz(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -823,7 +823,7 @@ namespace
 \r
     void xyz_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -849,7 +849,7 @@ namespace
 \r
     void xyz_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -875,7 +875,7 @@ namespace
 \r
     void rgb_to_hsv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -901,7 +901,7 @@ namespace
 \r
     void bgr_to_hsv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -927,7 +927,7 @@ namespace
 \r
     void hsv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -953,7 +953,7 @@ namespace
 \r
     void hsv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -979,7 +979,7 @@ namespace
 \r
     void rgb_to_hls(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1005,7 +1005,7 @@ namespace
 \r
     void bgr_to_hls(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1031,7 +1031,7 @@ namespace
 \r
     void hls_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1057,7 +1057,7 @@ namespace
 \r
     void hls_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1083,7 +1083,7 @@ namespace
 \r
     void rgb_to_hsv_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1109,7 +1109,7 @@ namespace
 \r
     void bgr_to_hsv_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1135,7 +1135,7 @@ namespace
 \r
     void hsv_to_rgb_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1161,7 +1161,7 @@ namespace
 \r
     void hsv_to_bgr_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1187,7 +1187,7 @@ namespace
 \r
     void rgb_to_hls_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1213,7 +1213,7 @@ namespace
 \r
     void bgr_to_hls_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1239,7 +1239,7 @@ namespace
 \r
     void hls_to_rgb_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1265,7 +1265,7 @@ namespace
 \r
     void hls_to_bgr_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
     {\r
-        using namespace cv::gpu::color;\r
+        using namespace cv::gpu::device;\r
         static const gpu_func_t funcs[2][2][6] = \r
         {\r
             {\r
@@ -1364,7 +1364,7 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream
         rgb_to_hls,             // CV_RGB2HLS     =53\r
 \r
         hsv_to_bgr,             // CV_HSV2BGR     =54\r
-        bgr_to_rgb,             // CV_HSV2RGB     =55\r
+        hsv_to_rgb,             // CV_HSV2RGB     =55\r
 \r
         0,                      // CV_Lab2BGR     =56\r
         0,                      // CV_Lab2RGB     =57\r

diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu
index 2755fff..a16c82c 100644 (file)
--- a/modules/gpu/src/cuda/color.cu
+++ b/modules/gpu/src/cuda/color.cu
@@ -44,11 +44,15 @@
 #include "opencv2/gpu/device/transform.hpp"\r
 #include "opencv2/gpu/device/color.hpp"\r
 \r
-using namespace cv::gpu;\r
-using namespace cv::gpu::device;\r
-\r
-namespace cv { namespace gpu { namespace color\r
+namespace cv { namespace gpu { namespace device\r
 {\r
+    template <> struct TransformFunctorTraits<bgra_to_rgba_traits<uchar>::functor_type> : DefaultTransformFunctorTraits<bgra_to_rgba_traits<uchar>::functor_type>\r
+    {\r
+        enum { smart_block_dim_x = 8 };\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
     #define OPENCV_GPU_IMPLEMENT_CVTCOLOR(name, traits) \\r
         void name(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream) \\r
         { \\r

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 (file)
--- 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
 \r
     namespace detail\r
     {\r
-        template <typename T, typename D, int bidx> struct RGB2RGB : public unary_function<T, D>\r
+        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>\r
         {\r
-            __device__ D operator()(const T& src) const\r
+            __device__ typename TypeVec<T, dcn>::vec_type operator()(const typename TypeVec<T, scn>::vec_type& src) const\r
             {\r
-                D dst;\r
+                typename TypeVec<T, dcn>::vec_type dst;\r
+\r
                 dst.x = (&src.x)[bidx];\r
                 dst.y = src.y;\r
                 dst.z = (&src.x)[bidx^2];\r
-                setAlpha(dst, getAlpha<typename VecTraits<T>::elem_type>(src));\r
+                setAlpha(dst, getAlpha<T>(src));\r
+\r
+                return dst;\r
+            }\r
+        };\r
+\r
+        template <> struct RGB2RGB<uchar, 4, 4, 2> : unary_function<uint, uint>\r
+        {\r
+            __device__ uint operator()(uint src) const\r
+            {\r
+                uint dst = 0;\r
+\r
+                dst |= (0xff & (src >> 16));\r
+                dst |= (0xff & (src >> 8)) << 8;\r
+                dst |= (0xff & (src)) << 16;\r
+                dst |= (0xff & (src >> 24)) << 24;\r
+\r
                 return dst;\r
             }\r
         };\r
@@ -117,10 +134,10 @@ namespace cv { namespace gpu { namespace device
 #define OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(name, scn, dcn, bidx) \\r
     template <typename T> struct name ## _traits \\r
     { \\r
-        typedef detail::RGB2RGB<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type, bidx> functor_type; \\r
+        typedef detail::RGB2RGB<T, scn, dcn, bidx> functor_type; \\r
         static __host__ __device__ __forceinline__ functor_type create_functor() \\r
         { \\r
-            return detail::RGB2RGB<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type, bidx>(); \\r
+            return functor_type(); \\r
         } \\r
     };\r
 \r

diff --git a/modules/gpu/test/test_imgproc.cpp b/modules/gpu/test/test_imgproc.cpp
index cc39a13..9f33783 100644 (file)
--- 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);\r
 }\r
 \r
+TEST_P(CvtColor, BGR2BGRA)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src = img;\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR2BGRA);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2BGRA);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
 TEST_P(CvtColor, BGRA2RGB)\r
 {\r
     PRINT_PARAM(devInfo);\r
     PRINT_TYPE(type);\r
 \r
     cv::Mat src;\r
-    cv::cvtColor(img, src, CV_BGR2BGRA);\r
+    cv::cvtColor(img, src, CV_BGR2BGRA);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGRA2RGB);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2RGB);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGRA2BGR)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGRA);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGRA2BGR);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2BGR);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGRA2RGBA)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGRA);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGRA2RGBA);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2RGBA);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR2GRAY)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src = img;\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR2GRAY);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2GRAY);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, RGB2GRAY)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2RGB);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_RGB2GRAY);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2GRAY);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, GRAY2BGR)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2GRAY);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_GRAY2BGR);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_GRAY2BGR);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, GRAY2BGRA)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2GRAY);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_GRAY2BGRA, 4);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_GRAY2BGRA, 4);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGRA2GRAY)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGRA);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGRA2GRAY);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2GRAY);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, RGBA2GRAY)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2RGBA);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_RGBA2GRAY);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGBA2GRAY);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, BGR2BGR565)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src = img;\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR2BGR565);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2BGR565);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, RGB2BGR565)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2RGB);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_RGB2BGR565);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2BGR565);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR5652BGR)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGR565);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR5652BGR);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5652BGR);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR5652RGB)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGR565);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR5652RGB);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5652RGB);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGRA2BGR565)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGRA);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGRA2BGR565);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2BGR565);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, RGBA2BGR565)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2RGBA);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_RGBA2BGR565);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGBA2BGR565);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR5652BGRA)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGR565);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR5652BGRA, 4);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5652BGRA, 4);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR5652RGBA)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGR565);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR5652RGBA, 4);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5652RGBA, 4);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, GRAY2BGR565)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2GRAY);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_GRAY2BGR565);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_GRAY2BGR565);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR5652GRAY)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGR565);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR5652GRAY);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5652GRAY);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR2BGR555)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src = img;\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR2BGR555);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2BGR555);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, RGB2BGR555)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2RGB);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_RGB2BGR555);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2BGR555);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR5552BGR)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGR555);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR5552BGR);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5552BGR);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR5552RGB)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGR555);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR5552RGB);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5552RGB);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGRA2BGR555)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGRA);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGRA2BGR555);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2BGR555);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, RGBA2BGR555)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2RGBA);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_RGBA2BGR555);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGBA2BGR555);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR5552BGRA)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGR555);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR5552BGRA, 4);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5552BGRA, 4);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR5552RGBA)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGR555);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR5552RGBA, 4);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5552RGBA, 4);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, GRAY2BGR555)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2GRAY);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_GRAY2BGR555);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_GRAY2BGR555);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR5552GRAY)\r
+{\r
+    if (type != CV_8U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2BGR555);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR5552GRAY);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR5552GRAY);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+}\r
+\r
+TEST_P(CvtColor, BGR2XYZ)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src = img;\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR2XYZ);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2XYZ);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, RGB2XYZ)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2RGB);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR2XYZ);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2XYZ);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, XYZ2BGR)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2XYZ);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_XYZ2BGR);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_XYZ2BGR);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, XYZ2RGB)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2XYZ);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_XYZ2RGB);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_XYZ2RGB);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, BGR2YCrCb)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src = img;\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_BGR2YCrCb);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2YCrCb);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, RGB2YCrCb)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2RGB);\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_BGRA2RGB);\r
+    cv::cvtColor(src, dst_gold, CV_RGB2YCrCb);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGRA2RGB);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2YCrCb);\r
 \r
         gpuRes.download(dst);\r
     );\r
 \r
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
 }\r
 \r
-TEST_P(CvtColor, BGR2YCrCb)\r
+TEST_P(CvtColor, YCrCb2BGR)\r
 {\r
     PRINT_PARAM(devInfo);\r
     PRINT_TYPE(type);\r
 \r
-    cv::Mat src = img;\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2YCrCb);\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_BGR2YCrCb);\r
+    cv::cvtColor(src, dst_gold, CV_YCrCb2BGR);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2YCrCb);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_YCrCb2BGR);\r
 \r
         gpuRes.download(dst);\r
     );\r
@@ -629,97 +1489,109 @@ TEST_P(CvtColor, YCrCb2RGB)
     EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
 }\r
 \r
-TEST_P(CvtColor, BGR2YUV)\r
+TEST_P(CvtColor, BGR2HSV)\r
 {\r
+    if (type == CV_16U)\r
+        return;\r
+\r
     PRINT_PARAM(devInfo);\r
     PRINT_TYPE(type);\r
 \r
     cv::Mat src = img;\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_BGR2YUV);\r
+    cv::cvtColor(src, dst_gold, CV_BGR2HSV);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2YUV);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HSV);\r
 \r
         gpuRes.download(dst);\r
     );\r
 \r
-    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
-TEST_P(CvtColor, YUV2BGR)\r
+TEST_P(CvtColor, RGB2HSV)\r
 {\r
+    if (type == CV_16U)\r
+        return;\r
+\r
     PRINT_PARAM(devInfo);\r
     PRINT_TYPE(type);\r
 \r
     cv::Mat src;\r
-    cv::cvtColor(img, src, CV_BGR2YUV);\r
+    cv::cvtColor(img, src, CV_BGR2RGB);\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_YUV2BGR);\r
+    cv::cvtColor(src, dst_gold, CV_RGB2HSV);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_YUV2BGR);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2HSV);\r
 \r
         gpuRes.download(dst);\r
     );\r
 \r
-    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
-TEST_P(CvtColor, BGR2XYZ)\r
+TEST_P(CvtColor, BGR2HLS)\r
 {\r
+    if (type == CV_16U)\r
+        return;\r
+\r
     PRINT_PARAM(devInfo);\r
     PRINT_TYPE(type);\r
 \r
     cv::Mat src = img;\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_BGR2XYZ);\r
+    cv::cvtColor(src, dst_gold, CV_BGR2HLS);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2XYZ);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HLS);\r
 \r
         gpuRes.download(dst);\r
     );\r
 \r
-    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
-TEST_P(CvtColor, XYZ2BGR)\r
+TEST_P(CvtColor, RGB2HLS)\r
 {\r
+    if (type == CV_16U)\r
+        return;\r
+\r
     PRINT_PARAM(devInfo);\r
     PRINT_TYPE(type);\r
 \r
     cv::Mat src;\r
-    cv::cvtColor(img, src, CV_BGR2XYZ);\r
+    cv::cvtColor(img, src, CV_BGR2RGB);\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_XYZ2BGR);\r
+    cv::cvtColor(src, dst_gold, CV_RGB2HLS);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_XYZ2BGR);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2HLS);\r
 \r
         gpuRes.download(dst);\r
     );\r
 \r
-    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
-TEST_P(CvtColor, BGR2HSV)\r
+TEST_P(CvtColor, HSV2BGR)\r
 {\r
     if (type == CV_16U)\r
         return;\r
@@ -727,16 +1599,17 @@ TEST_P(CvtColor, BGR2HSV)
     PRINT_PARAM(devInfo);\r
     PRINT_TYPE(type);\r
 \r
-    cv::Mat src = img;\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2HSV);\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_BGR2HSV);\r
+    cv::cvtColor(src, dst_gold, CV_HSV2BGR);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HSV);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2BGR);\r
 \r
         gpuRes.download(dst);\r
     );\r
@@ -744,7 +1617,7 @@ TEST_P(CvtColor, BGR2HSV)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
-TEST_P(CvtColor, HSV2BGR)\r
+TEST_P(CvtColor, HSV2RGB)\r
 {\r
     if (type == CV_16U)\r
         return;\r
@@ -755,14 +1628,14 @@ TEST_P(CvtColor, HSV2BGR)
     cv::Mat src;\r
     cv::cvtColor(img, src, CV_BGR2HSV);\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_HSV2BGR);\r
+    cv::cvtColor(src, dst_gold, CV_HSV2RGB);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2BGR);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2RGB);\r
 \r
         gpuRes.download(dst);\r
     );\r
@@ -770,7 +1643,7 @@ TEST_P(CvtColor, HSV2BGR)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
-TEST_P(CvtColor, BGR2HSV_FULL)\r
+TEST_P(CvtColor, HLS2BGR)\r
 {\r
     if (type == CV_16U)\r
         return;\r
@@ -778,16 +1651,17 @@ TEST_P(CvtColor, BGR2HSV_FULL)
     PRINT_PARAM(devInfo);\r
     PRINT_TYPE(type);\r
 \r
-    cv::Mat src = img;\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2HLS);\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_BGR2HSV_FULL);\r
+    cv::cvtColor(src, dst_gold, CV_HLS2BGR);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HSV_FULL);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HLS2BGR);\r
 \r
         gpuRes.download(dst);\r
     );\r
@@ -795,7 +1669,7 @@ TEST_P(CvtColor, BGR2HSV_FULL)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
-TEST_P(CvtColor, HSV2BGR_FULL)\r
+TEST_P(CvtColor, HLS2RGB)\r
 {\r
     if (type == CV_16U)\r
         return;\r
@@ -804,16 +1678,16 @@ TEST_P(CvtColor, HSV2BGR_FULL)
     PRINT_TYPE(type);\r
 \r
     cv::Mat src;\r
-    cv::cvtColor(img, src, CV_BGR2HSV_FULL);\r
+    cv::cvtColor(img, src, CV_BGR2HLS);\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_HSV2BGR_FULL);\r
+    cv::cvtColor(src, dst_gold, CV_HLS2RGB);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2BGR_FULL);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HLS2RGB);\r
 \r
         gpuRes.download(dst);\r
     );\r
@@ -821,7 +1695,7 @@ TEST_P(CvtColor, HSV2BGR_FULL)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
-TEST_P(CvtColor, BGR2HLS)\r
+TEST_P(CvtColor, BGR2HSV_FULL)\r
 {\r
     if (type == CV_16U)\r
         return;\r
@@ -831,14 +1705,14 @@ TEST_P(CvtColor, BGR2HLS)
 \r
     cv::Mat src = img;\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_BGR2HLS);\r
+    cv::cvtColor(src, dst_gold, CV_BGR2HSV_FULL);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HLS);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2HSV_FULL);\r
 \r
         gpuRes.download(dst);\r
     );\r
@@ -846,7 +1720,7 @@ TEST_P(CvtColor, BGR2HLS)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
-TEST_P(CvtColor, HLS2BGR)\r
+TEST_P(CvtColor, RGB2HSV_FULL)\r
 {\r
     if (type == CV_16U)\r
         return;\r
@@ -855,16 +1729,16 @@ TEST_P(CvtColor, HLS2BGR)
     PRINT_TYPE(type);\r
 \r
     cv::Mat src;\r
-    cv::cvtColor(img, src, CV_BGR2HLS);\r
+    cv::cvtColor(img, src, CV_BGR2RGB);\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_HLS2BGR);\r
+    cv::cvtColor(src, dst_gold, CV_RGB2HSV_FULL);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HLS2BGR);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2HSV_FULL);\r
 \r
         gpuRes.download(dst);\r
     );\r
@@ -897,6 +1771,84 @@ TEST_P(CvtColor, BGR2HLS_FULL)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
+TEST_P(CvtColor, RGB2HLS_FULL)\r
+{\r
+    if (type == CV_16U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2RGB);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_RGB2HLS_FULL);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2HLS_FULL);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
+}\r
+\r
+TEST_P(CvtColor, HSV2BGR_FULL)\r
+{\r
+    if (type == CV_16U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2HSV_FULL);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_HSV2BGR_FULL);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2BGR_FULL);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
+}\r
+\r
+TEST_P(CvtColor, HSV2RGB_FULL)\r
+{\r
+    if (type == CV_16U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2HSV_FULL);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_HSV2RGB_FULL);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HSV2RGB_FULL);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
+}\r
+\r
 TEST_P(CvtColor, HLS2BGR_FULL)\r
 {\r
     if (type == CV_16U)\r
@@ -923,21 +1875,47 @@ TEST_P(CvtColor, HLS2BGR_FULL)
     EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
 }\r
 \r
-TEST_P(CvtColor, BGR2GRAY)\r
+TEST_P(CvtColor, HLS2RGB_FULL)\r
+{\r
+    if (type == CV_16U)\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2HLS_FULL);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_HLS2RGB_FULL);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_HLS2RGB_FULL);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, type == CV_32F ? 1e-2 : 1);\r
+}\r
+\r
+TEST_P(CvtColor, BGR2YUV)\r
 {\r
     PRINT_PARAM(devInfo);\r
     PRINT_TYPE(type);\r
 \r
     cv::Mat src = img;\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_BGR2GRAY);\r
+    cv::cvtColor(src, dst_gold, CV_BGR2YUV);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2GRAY);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_BGR2YUV);\r
 \r
         gpuRes.download(dst);\r
     );\r
@@ -945,27 +1923,73 @@ TEST_P(CvtColor, BGR2GRAY)
     EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
 }\r
 \r
-TEST_P(CvtColor, GRAY2RGB)\r
+TEST_P(CvtColor, RGB2YUV)\r
 {\r
     PRINT_PARAM(devInfo);\r
     PRINT_TYPE(type);\r
 \r
     cv::Mat src;\r
-    cv::cvtColor(img, src, CV_BGR2GRAY);\r
+    cv::cvtColor(img, src, CV_BGR2RGB);\r
     cv::Mat dst_gold;\r
-    cv::cvtColor(src, dst_gold, CV_GRAY2RGB);\r
+    cv::cvtColor(src, dst_gold, CV_RGB2YUV);\r
 \r
     cv::Mat dst;\r
 \r
     ASSERT_NO_THROW(\r
         cv::gpu::GpuMat gpuRes;\r
 \r
-        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_GRAY2RGB);\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_RGB2YUV);\r
 \r
         gpuRes.download(dst);\r
     );\r
 \r
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, YUV2BGR)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_BGR2YUV);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_YUV2BGR);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_YUV2BGR);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
+}\r
+\r
+TEST_P(CvtColor, YUV2RGB)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+\r
+    cv::Mat src;\r
+    cv::cvtColor(img, src, CV_RGB2YUV);\r
+    cv::Mat dst_gold;\r
+    cv::cvtColor(src, dst_gold, CV_YUV2RGB);\r
+\r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpuRes;\r
+\r
+        cv::gpu::cvtColor(cv::gpu::GpuMat(src), gpuRes, CV_YUV2RGB);\r
+\r
+        gpuRes.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor, testing::Combine(\r