fixed gpu::cvtColor according #5324 changes

diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp
index e95ce86..4d9bd50 100644 (file)
--- a/modules/gpu/src/color.cpp
+++ b/modules/gpu/src/color.cpp
@@ -411,7 +411,7 @@ namespace
                     bidx = code == CV_BGR2HSV || code == CV_BGR2HLS ||\r
                         code == CV_BGR2HSV_FULL || code == CV_BGR2HLS_FULL ? 0 : 2;\r
                     int hrange = depth == CV_32F ? 360 : code == CV_BGR2HSV || code == CV_RGB2HSV ||\r
-                        code == CV_BGR2HLS || code == CV_RGB2HLS ? 180 : 255;\r
+                        code == CV_BGR2HLS || code == CV_RGB2HLS ? 180 : 256;\r
                 \r
                     dst.create(sz, CV_MAKETYPE(depth, dcn));\r
 \r

diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu
index e3d9851..7a0c45a 100644 (file)
--- a/modules/gpu/src/cuda/color.cu
+++ b/modules/gpu/src/cuda/color.cu
@@ -913,47 +913,15 @@ namespace cv { namespace gpu { namespace color
 \r
 ////////////////////////////////////// RGB <-> HSV ///////////////////////////////////////\r
 \r
-    __constant__ int cHsvDivTable[256] = \r
-    {\r
-        0, 1044480, 522240, 348160, 261120, 208896, 174080, 149211,\r
-        130560, 116053, 104448, 94953, 87040, 80345, 74606, 69632,\r
-        65280, 61440, 58027, 54973, 52224, 49737, 47476, 45412,\r
-        43520, 41779, 40172, 38684, 37303, 36017, 34816, 33693,\r
-        32640, 31651, 30720, 29842, 29013, 28229, 27486, 26782,\r
-        26112, 25475, 24869, 24290, 23738, 23211, 22706, 22223,\r
-        21760, 21316, 20890, 20480, 20086, 19707, 19342, 18991,\r
-        18651, 18324, 18008, 17703, 17408, 17123, 16846, 16579,\r
-        16320, 16069, 15825, 15589, 15360, 15137, 14921, 14711,\r
-        14507, 14308, 14115, 13926, 13743, 13565, 13391, 13221,\r
-        13056, 12895, 12738, 12584, 12434, 12288, 12145, 12006,\r
-        11869, 11736, 11605, 11478, 11353, 11231, 11111, 10995,\r
-        10880, 10768, 10658, 10550, 10445, 10341, 10240, 10141,\r
-        10043, 9947, 9854, 9761, 9671, 9582, 9495, 9410,\r
-        9326, 9243, 9162, 9082, 9004, 8927, 8852, 8777,\r
-        8704, 8632, 8561, 8492, 8423, 8356, 8290, 8224,\r
-        8160, 8097, 8034, 7973, 7913, 7853, 7795, 7737,\r
-        7680, 7624, 7569, 7514, 7461, 7408, 7355, 7304,\r
-        7253, 7203, 7154, 7105, 7057, 7010, 6963, 6917,\r
-        6872, 6827, 6782, 6739, 6695, 6653, 6611, 6569,\r
-        6528, 6487, 6447, 6408, 6369, 6330, 6292, 6254,\r
-        6217, 6180, 6144, 6108, 6073, 6037, 6003, 5968,\r
-        5935, 5901, 5868, 5835, 5803, 5771, 5739, 5708,\r
-        5677, 5646, 5615, 5585, 5556, 5526, 5497, 5468,\r
-        5440, 5412, 5384, 5356, 5329, 5302, 5275, 5249,\r
-        5222, 5196, 5171, 5145, 5120, 5095, 5070, 5046,\r
-        5022, 4998, 4974, 4950, 4927, 4904, 4881, 4858,\r
-        4836, 4813, 4791, 4769, 4748, 4726, 4705, 4684,\r
-        4663, 4642, 4622, 4601, 4581, 4561, 4541, 4522,\r
-        4502, 4483, 4464, 4445, 4426, 4407, 4389, 4370,\r
-        4352, 4334, 4316, 4298, 4281, 4263, 4246, 4229,\r
-        4212, 4195, 4178, 4161, 4145, 4128, 4112, 4096\r
-    };\r
+    __constant__ int cHsvDivTable   [256] = {0, 1044480, 522240, 348160, 261120, 208896, 174080, 149211, 130560, 116053, 104448, 94953, 87040, 80345, 74606, 69632, 65280, 61440, 58027, 54973, 52224, 49737, 47476, 45412, 43520, 41779, 40172, 38684, 37303, 36017, 34816, 33693, 32640, 31651, 30720, 29842, 29013, 28229, 27486, 26782, 26112, 25475, 24869, 24290, 23738, 23211, 22706, 22223, 21760, 21316, 20890, 20480, 20086, 19707, 19342, 18991, 18651, 18324, 18008, 17703, 17408, 17123, 16846, 16579, 16320, 16069, 15825, 15589, 15360, 15137, 14921, 14711, 14507, 14308, 14115, 13926, 13743, 13565, 13391, 13221, 13056, 12895, 12738, 12584, 12434, 12288, 12145, 12006, 11869, 11736, 11605, 11478, 11353, 11231, 11111, 10995, 10880, 10768, 10658, 10550, 10445, 10341, 10240, 10141, 10043, 9947, 9854, 9761, 9671, 9582, 9495, 9410, 9326, 9243, 9162, 9082, 9004, 8927, 8852, 8777, 8704, 8632, 8561, 8492, 8423, 8356, 8290, 8224, 8160, 8097, 8034, 7973, 7913, 7853, 7795, 7737, 7680, 7624, 7569, 7514, 7461, 7408, 7355, 7304, 7253, 7203, 7154, 7105, 7057, 7010, 6963, 6917, 6872, 6827, 6782, 6739, 6695, 6653, 6611, 6569, 6528, 6487, 6447, 6408, 6369, 6330, 6292, 6254, 6217, 6180, 6144, 6108, 6073, 6037, 6003, 5968, 5935, 5901, 5868, 5835, 5803, 5771, 5739, 5708, 5677, 5646, 5615, 5585, 5556, 5526, 5497, 5468, 5440, 5412, 5384, 5356, 5329, 5302, 5275, 5249, 5222, 5196, 5171, 5145, 5120, 5095, 5070, 5046, 5022, 4998, 4974, 4950, 4927, 4904, 4881, 4858, 4836, 4813, 4791, 4769, 4748, 4726, 4705, 4684, 4663, 4642, 4622, 4601, 4581, 4561, 4541, 4522, 4502, 4483, 4464, 4445, 4426, 4407, 4389, 4370, 4352, 4334, 4316, 4298, 4281, 4263, 4246, 4229, 4212, 4195, 4178, 4161, 4145, 4128, 4112, 4096};\r
+    __constant__ int cHsvDivTable180[256] = {0, 122880, 61440, 40960, 30720, 24576, 20480, 17554, 15360, 13653, 12288, 11171, 10240, 9452, 8777, 8192, 7680, 7228, 6827, 6467, 6144, 5851, 5585, 5343, 5120, 4915, 4726, 4551, 4389, 4237, 4096, 3964, 3840, 3724, 3614, 3511, 3413, 3321, 3234, 3151, 3072, 2997, 2926, 2858, 2793, 2731, 2671, 2614, 2560, 2508, 2458, 2409, 2363, 2318, 2276, 2234, 2194, 2156, 2119, 2083, 2048, 2014, 1982, 1950, 1920, 1890, 1862, 1834, 1807, 1781, 1755, 1731, 1707, 1683, 1661, 1638, 1617, 1596, 1575, 1555, 1536, 1517, 1499, 1480, 1463, 1446, 1429, 1412, 1396, 1381, 1365, 1350, 1336, 1321, 1307, 1293, 1280, 1267, 1254, 1241, 1229, 1217, 1205, 1193, 1182, 1170, 1159, 1148, 1138, 1127, 1117, 1107, 1097, 1087, 1078, 1069, 1059, 1050, 1041, 1033, 1024, 1016, 1007, 999, 991, 983, 975, 968, 960, 953, 945, 938, 931, 924, 917, 910, 904, 897, 890, 884, 878, 871, 865, 859, 853, 847, 842, 836, 830, 825, 819, 814, 808, 803, 798, 793, 788, 783, 778, 773, 768, 763, 759, 754, 749, 745, 740, 736, 731, 727, 723, 719, 714, 710, 706, 702, 698, 694, 690, 686, 683, 679, 675, 671, 668, 664, 661, 657, 654, 650, 647, 643, 640, 637, 633, 630, 627, 624, 621, 617, 614, 611, 608, 605, 602, 599, 597, 594, 591, 588, 585, 582, 580, 577, 574, 572, 569, 566, 564, 561, 559, 556, 554, 551, 549, 546, 544, 541, 539, 537, 534, 532, 530, 527, 525, 523, 521, 518, 516, 514, 512, 510, 508, 506, 504, 502, 500, 497, 495, 493, 492, 490, 488, 486, 484, 482};\r
+    __constant__ int cHsvDivTable256[256] = {0, 174763, 87381, 58254, 43691, 34953, 29127, 24966, 21845, 19418, 17476, 15888, 14564, 13443, 12483, 11651, 10923, 10280, 9709, 9198, 8738, 8322, 7944, 7598, 7282, 6991, 6722, 6473, 6242, 6026, 5825, 5638, 5461, 5296, 5140, 4993, 4855, 4723, 4599, 4481, 4369, 4263, 4161, 4064, 3972, 3884, 3799, 3718, 3641, 3567, 3495, 3427, 3361, 3297, 3236, 3178, 3121, 3066, 3013, 2962, 2913, 2865, 2819, 2774, 2731, 2689, 2648, 2608, 2570, 2533, 2497, 2461, 2427, 2394, 2362, 2330, 2300, 2270, 2241, 2212, 2185, 2158, 2131, 2106, 2081, 2056, 2032, 2009, 1986, 1964, 1942, 1920, 1900, 1879, 1859, 1840, 1820, 1802, 1783, 1765, 1748, 1730, 1713, 1697, 1680, 1664, 1649, 1633, 1618, 1603, 1589, 1574, 1560, 1547, 1533, 1520, 1507, 1494, 1481, 1469, 1456, 1444, 1432, 1421, 1409, 1398, 1387, 1376, 1365, 1355, 1344, 1334, 1324, 1314, 1304, 1295, 1285, 1276, 1266, 1257, 1248, 1239, 1231, 1222, 1214, 1205, 1197, 1189, 1181, 1173, 1165, 1157, 1150, 1142, 1135, 1128, 1120, 1113, 1106, 1099, 1092, 1085, 1079, 1072, 1066, 1059, 1053, 1046, 1040, 1034, 1028, 1022, 1016, 1010, 1004, 999, 993, 987, 982, 976, 971, 966, 960, 955, 950, 945, 940, 935, 930, 925, 920, 915, 910, 906, 901, 896, 892, 887, 883, 878, 874, 869, 865, 861, 857, 853, 848, 844, 840, 836, 832, 828, 824, 820, 817, 813, 809, 805, 802, 798, 794, 791, 787, 784, 780, 777, 773, 770, 767, 763, 760, 757, 753, 750, 747, 744, 741, 737, 734, 731, 728, 725, 722, 719, 716, 713, 710, 708, 705, 702, 699, 696, 694, 691, 688, 685};\r
     \r
     template <int HR, typename D>\r
     __device__ void RGB2HSVConvert(const uchar* src, D& dst, int bidx)\r
     {\r
         const int hsv_shift = 12;\r
-        const int hscale = HR == 180 ? 15 : 21;\r
+        const int* hdiv_table = HR == 180 ? cHsvDivTable180 : cHsvDivTable256;\r
 \r
         int b = src[bidx], g = src[1], r = src[bidx^2];\r
         int h, s, v = b;\r
@@ -969,12 +937,12 @@ namespace cv { namespace gpu { namespace color
         vr = v == r ? -1 : 0;\r
         vg = v == g ? -1 : 0;\r
 \r
-        s = diff * cHsvDivTable[v] >> hsv_shift;\r
+        s = (diff * cHsvDivTable[v] + (1 << (hsv_shift-1))) >> hsv_shift;\r
         h = (vr & (g - b)) + (~vr & ((vg & (b - r + 2 * diff)) + ((~vg) & (r - g + 4 * diff))));\r
-        h = (h * cHsvDivTable[diff] * hscale + (1 << (hsv_shift + 6))) >> (7 + hsv_shift);\r
+        h = (h * hdiv_table[diff] + (1 << (hsv_shift-1))) >> hsv_shift;\r
         h += h < 0 ? HR : 0;\r
 \r
-        dst.x = (uchar)h;\r
+        dst.x = saturate_cast<uchar>(h);\r
         dst.y = (uchar)s;\r
         dst.z = (uchar)v;\r
     }\r
@@ -1116,7 +1084,7 @@ namespace cv { namespace gpu { namespace color
         }\r
         else\r
         {\r
-            RGB2HSV<SRCCN, DSTCN, 255, T> cvt(bidx);\r
+            RGB2HSV<SRCCN, DSTCN, 256, T> cvt(bidx);\r
             callConvert(src, dst, cvt, stream);\r
         }\r
     }\r
@@ -1349,7 +1317,7 @@ namespace cv { namespace gpu { namespace color
         }\r
         else\r
         {\r
-            RGB2HLS<SRCCN, DSTCN, 255, T> cvt(bidx);\r
+            RGB2HLS<SRCCN, DSTCN, 256, T> cvt(bidx);\r
             callConvert(src, dst, cvt, stream);\r
         }\r
     }\r