Skip to content

O
opencv
Commit 0d09352f authored by Vladislav Vinogradov's avatar Vladislav Vinogradov
Browse files
Options

fixed gpu::cvtColor

parent ada3e6e6
Show whitespace changes
Inline Side-by-side
Showing with 49 additions and 92 deletions
modules/gpu/src/cuda/color.cu
View file @ 0d09352f
......@@ -55,18 +55,11 @@ using namespace cv::gpu::device;
namespace cv { namespace gpu { namespace color
{
template<typename T> struct ColorChannel;
template<> struct ColorChannel<uchar>
template<typename T> struct ColorChannel
{
typedef float worktype_f;
static __device__ __forceinline__ uchar max() { return UCHAR_MAX; }
static __device__ __forceinline__ uchar half() { return (uchar)(max()/2 + 1); }
};
template<> struct ColorChannel<ushort>
{
typedef float worktype_f;
static __device__ __forceinline__ ushort max() { return USHRT_MAX; }
static __device__ __forceinline__ ushort half() { return (ushort)(max()/2 + 1); }
static __device__ __forceinline__ T max() { return numeric_limits_gpu<T>::max(); }
static __device__ __forceinline__ T half() { return (T)(max()/2 + 1); }
};
template<> struct ColorChannel<float>
{
......@@ -917,11 +910,10 @@ namespace cv { namespace gpu { namespace color
__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};
__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};
template <int HR, typename D>
__device__ void RGB2HSVConvert(const uchar* src, D& dst, int bidx)
template <typename D> __device__ void RGB2HSVConvert(const uchar* src, D& dst, int bidx, int hr)
{
const int hsv_shift = 12;
const int* hdiv_table = HR == 180 ? cHsvDivTable180 : cHsvDivTable256;
const int* hdiv_table = hr == 180 ? cHsvDivTable180 : cHsvDivTable256;
int b = src[bidx], g = src[1], r = src[bidx^2];
int h, s, v = b;
......@@ -940,16 +932,15 @@ namespace cv { namespace gpu { namespace color
s = (diff * cHsvDivTable[v] + (1 << (hsv_shift-1))) >> hsv_shift;
h = (vr & (g - b)) + (~vr & ((vg & (b - r + 2 * diff)) + ((~vg) & (r - g + 4 * diff))));
h = (h * hdiv_table[diff] + (1 << (hsv_shift-1))) >> hsv_shift;
h += h < 0 ? HR : 0;
h += h < 0 ? hr : 0;
dst.x = saturate_cast<uchar>(h);
dst.y = (uchar)s;
dst.z = (uchar)v;
}
template<int HR, typename D>
__device__ void RGB2HSVConvert(const float* src, D& dst, int bidx)
template<typename D> __device__ void RGB2HSVConvert(const float* src, D& dst, int bidx, int hr)
{
const float hscale = HR * (1.f / 360.f);
const float hscale = hr * (1.f / 360.f);
float b = src[bidx], g = src[1], r = src[bidx^2];
float h, s, v;
......@@ -980,22 +971,23 @@ namespace cv { namespace gpu { namespace color
dst.z = v;
}
template <int SRCCN, int DSTCN, int HR, typename T> struct RGB2HSV
template <int SRCCN, int DSTCN, typename T> struct RGB2HSV
{
typedef typename TypeVec<T, SRCCN>::vec_t src_t;
typedef typename TypeVec<T, DSTCN>::vec_t dst_t;
explicit RGB2HSV(int bidx) : bidx(bidx) {}
RGB2HSV(int bidx, int hr) : bidx(bidx), hr(hr) {}
__device__ __forceinline__ dst_t operator()(const src_t& src) const
{
dst_t dst;
RGB2HSVConvert<HR>(&src.x, dst, bidx);
RGB2HSVConvert(&src.x, dst, bidx, hr);
return dst;
}
private:
int bidx;
int hr;
};
__constant__ int cHsvSectorData[6][3] =
......@@ -1003,10 +995,9 @@ namespace cv { namespace gpu { namespace color
{1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0}
};
template <int HR, typename T>
__device__ void HSV2RGBConvert(const T& src, float* dst, int bidx)
template <typename T> __device__ void HSV2RGBConvert(const T& src, float* dst, int bidx, int hr)
{
const float hscale = 6.f / HR;
const float hscale = 6.f / hr;
float h = src.x, s = src.y, v = src.z;
float b, g, r;
......@@ -1039,8 +1030,7 @@ namespace cv { namespace gpu { namespace color
dst[1] = g;
dst[bidx^2] = r;
}
template <int HR, typename T>
__device__ void HSV2RGBConvert(const T& src, uchar* dst, int bidx)
template <typename T> __device__ void HSV2RGBConvert(const T& src, uchar* dst, int bidx, int hr)
{
float3 buf;
......@@ -1048,46 +1038,39 @@ namespace cv { namespace gpu { namespace color
buf.y = src.y * (1.f/255.f);
buf.z = src.z * (1.f/255.f);
HSV2RGBConvert<HR>(buf, &buf.x, bidx);
HSV2RGBConvert(buf, &buf.x, bidx, hr);
dst[0] = saturate_cast<uchar>(buf.x * 255.f);
dst[1] = saturate_cast<uchar>(buf.y * 255.f);
dst[2] = saturate_cast<uchar>(buf.z * 255.f);
}
template <int SRCCN, int DSTCN, int HR, typename T> struct HSV2RGB
template <int SRCCN, int DSTCN, typename T> struct HSV2RGB
{
typedef typename TypeVec<T, SRCCN>::vec_t src_t;
typedef typename TypeVec<T, DSTCN>::vec_t dst_t;
explicit HSV2RGB(int bidx) : bidx(bidx) {}
HSV2RGB(int bidx, int hr) : bidx(bidx), hr(hr) {}
__device__ __forceinline__ dst_t operator()(const src_t& src) const
{
dst_t dst;
HSV2RGBConvert<HR>(src, &dst.x, bidx);
HSV2RGBConvert(src, &dst.x, bidx, hr);
setAlpha(dst, ColorChannel<T>::max());
return dst;
}
private:
int bidx;
int hr;
};
template <typename T, int SRCCN, int DSTCN>
void RGB2HSV_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream)
{
if (hrange == 180)
{
RGB2HSV<SRCCN, DSTCN, 180, T> cvt(bidx);
callConvert(src, dst, cvt, stream);
}
else
{
RGB2HSV<SRCCN, DSTCN, 256, T> cvt(bidx);
RGB2HSV<SRCCN, DSTCN, T> cvt(bidx, hrange);
callConvert(src, dst, cvt, stream);
}
}
void RGB2HSV_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream)
{
......@@ -1116,17 +1099,9 @@ namespace cv { namespace gpu { namespace color
template <typename T, int SRCCN, int DSTCN>
void HSV2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream)
{
if (hrange == 180)
{
HSV2RGB<SRCCN, DSTCN, 180, T> cvt(bidx);
callConvert(src, dst, cvt, stream);
}
else
{
HSV2RGB<SRCCN, DSTCN, 255, T> cvt(bidx);
HSV2RGB<SRCCN, DSTCN, T> cvt(bidx, hrange);
callConvert(src, dst, cvt, stream);
}
}
void HSV2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream)
{
......@@ -1154,10 +1129,9 @@ namespace cv { namespace gpu { namespace color
/////////////////////////////////////// RGB <-> HLS ////////////////////////////////////////
template <int HR, typename D>
__device__ void RGB2HLSConvert(const float* src, D& dst, int bidx)
template <typename D> __device__ void RGB2HLSConvert(const float* src, D& dst, int bidx, int hr)
{
const float hscale = HR * (1.f/360.f);
const float hscale = hr * (1.f / 360.f);
float b = src[bidx], g = src[1], r = src[bidx^2];
float h = 0.f, s = 0.f, l;
......@@ -1191,8 +1165,7 @@ namespace cv { namespace gpu { namespace color
dst.y = l;
dst.z = s;
}
template <int HR, typename D>
__device__ void RGB2HLSConvert(const uchar* src, D& dst, int bidx)
template <typename D> __device__ void RGB2HLSConvert(const uchar* src, D& dst, int bidx, int hr)
{
float3 buf;
......@@ -1200,29 +1173,30 @@ namespace cv { namespace gpu { namespace color
buf.y = src[1]*(1.f/255.f);
buf.z = src[2]*(1.f/255.f);
RGB2HLSConvert<HR>(&buf.x, buf, bidx);
RGB2HLSConvert(&buf.x, buf, bidx, hr);
dst.x = saturate_cast<uchar>(buf.x);
dst.y = saturate_cast<uchar>(buf.y*255.f);
dst.z = saturate_cast<uchar>(buf.z*255.f);
}
template <int SRCCN, int DSTCN, int HR, typename T> struct RGB2HLS
template <int SRCCN, int DSTCN, typename T> struct RGB2HLS
{
typedef typename TypeVec<T, SRCCN>::vec_t src_t;
typedef typename TypeVec<T, DSTCN>::vec_t dst_t;
explicit RGB2HLS(int bidx) : bidx(bidx) {}
RGB2HLS(int bidx, int hr) : bidx(bidx), hr(hr) {}
__device__ __forceinline__ dst_t operator()(const src_t& src) const
{
dst_t dst;
RGB2HLSConvert<HR>(&src.x, dst, bidx);
RGB2HLSConvert(&src.x, dst, bidx, hr);
return dst;
}
private:
int bidx;
int hr;
};
__constant__ int cHlsSectorData[6][3] =
......@@ -1230,10 +1204,9 @@ namespace cv { namespace gpu { namespace color
{1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0}
};
template <int HR, typename T>
__device__ void HLS2RGBConvert(const T& src, float* dst, int bidx)
template <typename T> __device__ void HLS2RGBConvert(const T& src, float* dst, int bidx, int hr)
{
const float hscale = 6.0f / HR;
const float hscale = 6.0f / hr;
float h = src.x, l = src.y, s = src.z;
float b, g, r;
......@@ -1272,8 +1245,7 @@ namespace cv { namespace gpu { namespace color
dst[1] = g;
dst[bidx^2] = r;
}
template <int HR, typename T>
__device__ void HLS2RGBConvert(const T& src, uchar* dst, int bidx)
template <typename T> __device__ void HLS2RGBConvert(const T& src, uchar* dst, int bidx, int hr)
{
float3 buf;
......@@ -1281,46 +1253,39 @@ namespace cv { namespace gpu { namespace color
buf.y = src.y*(1.f/255.f);
buf.z = src.z*(1.f/255.f);
HLS2RGBConvert<HR>(buf, &buf.x, bidx);
HLS2RGBConvert(buf, &buf.x, bidx, hr);
dst[0] = saturate_cast<uchar>(buf.x*255.f);
dst[1] = saturate_cast<uchar>(buf.y*255.f);
dst[2] = saturate_cast<uchar>(buf.z*255.f);
}
template <int SRCCN, int DSTCN, int HR, typename T> struct HLS2RGB
template <int SRCCN, int DSTCN, typename T> struct HLS2RGB
{
typedef typename TypeVec<T, SRCCN>::vec_t src_t;
typedef typename TypeVec<T, DSTCN>::vec_t dst_t;
explicit HLS2RGB(int bidx) : bidx(bidx) {}
HLS2RGB(int bidx, int hr) : bidx(bidx), hr(hr) {}
__device__ __forceinline__ dst_t operator()(const src_t& src) const
{
dst_t dst;
HLS2RGBConvert<HR>(src, &dst.x, bidx);
HLS2RGBConvert(src, &dst.x, bidx, hr);
setAlpha(dst, ColorChannel<T>::max());
return dst;
}
private:
int bidx;
int hr;
};
template <typename T, int SRCCN, int DSTCN>
void RGB2HLS_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream)
{
if (hrange == 180)
{
RGB2HLS<SRCCN, DSTCN, 180, T> cvt(bidx);
RGB2HLS<SRCCN, DSTCN, T> cvt(bidx, hrange);
callConvert(src, dst, cvt, stream);
}
else
{
RGB2HLS<SRCCN, DSTCN, 256, T> cvt(bidx);
callConvert(src, dst, cvt, stream);
}
}
void RGB2HLS_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream)
{
......@@ -1350,17 +1315,9 @@ namespace cv { namespace gpu { namespace color
template <typename T, int SRCCN, int DSTCN>
void HLS2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream)
{
if (hrange == 180)
{
HLS2RGB<SRCCN, DSTCN, 180, T> cvt(bidx);
callConvert(src, dst, cvt, stream);
}
else
{
HLS2RGB<SRCCN, DSTCN, 255, T> cvt(bidx);
HLS2RGB<SRCCN, DSTCN, T> cvt(bidx, hrange);
callConvert(src, dst, cvt, stream);
}
}
void HLS2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream)
{
......
modules/gpu/test/test_imgproc_gpu.cpp
View file @ 0d09352f
......@@ -461,17 +461,17 @@ protected:
int CV_GpuCvtColorTest::CheckNorm(const Mat& m1, const Mat& m2)
{
double ret = norm(m1, m2, NORM_INF);
float max_err = 1e-2f;
Mat diff;
cv::matchTemplate(m1, m2, diff, CV_TM_CCORR_NORMED);
float err = abs(diff.at<float>(0, 0) - 1.f);
if (err > max_err)
return cvtest::TS::FAIL_INVALID_OUTPUT;
if (ret <= 3)
{
return cvtest::TS::OK;
}
else
{
ts->printf(cvtest::TS::LOG, "\nNorm: %f\n", ret);
return cvtest::TS::FAIL_GENERIC;
}
}
void CV_GpuCvtColorTest::run( int )
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment