/*
case COLOR_BGR2BGR565: case COLOR_BGR2BGR555: case COLOR_RGB2BGR565: case COLOR_RGB2BGR555:
case COLOR_BGRA2BGR565: case COLOR_BGRA2BGR555: case COLOR_RGBA2BGR565: case COLOR_RGBA2BGR555:
- case COLOR_BGR5652BGR: case COLOR_BGR5552BGR: case COLOR_BGR5652RGB: case COLOR_BGR5552RGB:
- case COLOR_BGR5652BGRA: case COLOR_BGR5552BGRA: case COLOR_BGR5652RGBA: case COLOR_BGR5552RGBA:
*/
case COLOR_BGR2BGRA: case COLOR_RGB2BGRA: case COLOR_BGRA2BGR:
case COLOR_RGBA2BGR: case COLOR_RGB2BGR: case COLOR_BGRA2RGBA:
reverse ? "REVERSE" : "ORDER"));
break;
}
+ case COLOR_BGR5652BGR: case COLOR_BGR5552BGR: case COLOR_BGR5652RGB: case COLOR_BGR5552RGB:
+ case COLOR_BGR5652BGRA: case COLOR_BGR5552BGRA: case COLOR_BGR5652RGBA: case COLOR_BGR5552RGBA:
+ {
+ dcn = code == COLOR_BGR5652BGRA || code == COLOR_BGR5552BGRA || code == COLOR_BGR5652RGBA || code == COLOR_BGR5552RGBA ? 4 : 3;
+ CV_Assert((dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U);
+ bidx = code == COLOR_BGR5652BGR || code == COLOR_BGR5552BGR ||
+ code == COLOR_BGR5652BGRA || code == COLOR_BGR5552BGRA ? 0 : 2;
+ int greenbits = code == COLOR_BGR5652BGR || code == COLOR_BGR5652RGB ||
+ code == COLOR_BGR5652BGRA || code == COLOR_BGR5652RGBA ? 6 : 5;
+ k.create("RGB5x52RGB", ocl::imgproc::cvtcolor_oclsrc,
+ format("-D depth=%d -D scn=2 -D dcn=%s -D bidx=%d -D greenbits=%d", depth, scn, dcn, bidx, greenbits));
+ break;
+ }
case COLOR_BGR2GRAY: case COLOR_BGRA2GRAY:
case COLOR_RGB2GRAY: case COLOR_RGBA2GRAY:
{
}
+///////////////////////////////////// RGB5x5 <-> RGB //////////////////////////////////////
+
+__kernel void RGB5x52RGB(__global const uchar* src, int src_step, int src_offset,
+ __global uchar* dst, int dst_step, int dst_offset,
+ int rows, int cols)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
+
+ if (y < rows && x < cols)
+ {
+ int src_idx = mad24(y, src_step, src_offset + x * scnbytes);
+ int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes);
+ ushort t = *((__global const ushort*)(src + src_idx));
+
+#if greenbits == 6
+ dst[dst_idx + bidx] = (uchar)(t << 3);
+ dst[dst_idx + 1] = (uchar)((t >> 3) & ~3);
+ dst[dst_idx + (bidx^2)] = (uchar)((t >> 8) & ~7);
+#else
+ dst[dst_idx + bidx] = (uchar)(t << 3);
+ dst[dst_idx + 1] = (uchar)((t >> 2) & ~7);
+ dst[dst_idx + (bidx^2)] = (uchar)((t >> 7) & ~7);
+#endif
+
+#if dcn == 4
+#if greenbits == 6
+ dst[dst_idx + 3] = 255;
+#else
+ dst[dst_idx + 3] = t & 0x8000 ? 255 : 0;
+#endif
+#endif
+ }
+}
+
+//__kernel void RGB2RGB5x5(int cols, int rows, int src_step, int dst_step, int bidx,
+// __global const uchar * src, __global ushort * dst,
+// int src_offset, int dst_offset)
+//{
+// int x = get_global_id(0);
+// int y = get_global_id(1);
+
+// if (y < rows && x < cols)
+// {
+// int src_idx = mad24(y, src_step, src_offset + (x << 2));
+// int dst_idx = mad24(y, dst_step, dst_offset + x);
+
+//#if greenbits == 6
+// dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~3) << 3)|((src[src_idx + (bidx^2)]&~7) << 8));
+//#elif scn == 3
+// dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)|((src[src_idx + (bidx^2)]&~7) << 7));
+//#else
+// dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)|
+// ((src[src_idx + (bidx^2)]&~7) << 7)|(src[src_idx + 3] ? 0x8000 : 0));
+//#endif
+// }
+//}
// RGB5x5 <-> RGB
-//typedef CvtColor CvtColor8u;
+typedef CvtColor CvtColor8u;
-//OCL_TEST_P(CvtColor8u, BGR5652BGR) { performTest(2, 3, CVTCODE(BGR5652BGR)); }
-//OCL_TEST_P(CvtColor8u, BGR5652RGB) { performTest(2, 3, CVTCODE(BGR5652RGB)); }
-//OCL_TEST_P(CvtColor8u, BGR5652BGRA) { performTest(2, 4, CVTCODE(BGR5652BGRA)); }
-//OCL_TEST_P(CvtColor8u, BGR5652RGBA) { performTest(2, 4, CVTCODE(BGR5652RGBA)); }
+OCL_TEST_P(CvtColor8u, BGR5652BGR) { performTest(2, 3, CVTCODE(BGR5652BGR)); }
+OCL_TEST_P(CvtColor8u, BGR5652RGB) { performTest(2, 3, CVTCODE(BGR5652RGB)); }
+OCL_TEST_P(CvtColor8u, BGR5652BGRA) { performTest(2, 4, CVTCODE(BGR5652BGRA)); }
+OCL_TEST_P(CvtColor8u, BGR5652RGBA) { performTest(2, 4, CVTCODE(BGR5652RGBA)); }
-//OCL_TEST_P(CvtColor8u, BGR5552BGR) { performTest(2, 3, CVTCODE(BGR5552BGR)); }
-//OCL_TEST_P(CvtColor8u, BGR5552RGB) { performTest(2, 3, CVTCODE(BGR5552RGB)); }
-//OCL_TEST_P(CvtColor8u, BGR5552BGRA) { performTest(2, 4, CVTCODE(BGR5552BGRA)); }
-//OCL_TEST_P(CvtColor8u, BGR5552RGBA) { performTest(2, 4, CVTCODE(BGR5552RGBA)); }
+OCL_TEST_P(CvtColor8u, BGR5552BGR) { performTest(2, 3, CVTCODE(BGR5552BGR)); }
+OCL_TEST_P(CvtColor8u, BGR5552RGB) { performTest(2, 3, CVTCODE(BGR5552RGB)); }
+OCL_TEST_P(CvtColor8u, BGR5552BGRA) { performTest(2, 4, CVTCODE(BGR5552BGRA)); }
+OCL_TEST_P(CvtColor8u, BGR5552RGBA) { performTest(2, 4, CVTCODE(BGR5552RGBA)); }
//OCL_TEST_P(CvtColor8u, BGR2BGR565) { performTest(3, 2, CVTCODE(BGR2BGR565)); }
//OCL_TEST_P(CvtColor8u, RGB2BGR565) { performTest(3, 2, CVTCODE(RGB2BGR565)); }
BLOCK_SIZE = 256
};
-///////////////////////////////////// RGB[A] <-> BGR[A] //////////////////////////////////////
-
-__kernel void RGB(int cols, int rows, int src_step, int dst_step,
- __global const DATA_TYPE * src, __global DATA_TYPE * dst,
- int src_offset, int dst_offset)
-{
- int x = get_global_id(0);
- int y = get_global_id(1);
-
- if (y < rows && x < cols)
- {
- x <<= 2;
- int src_idx = mad24(y, src_step, src_offset + x);
- int dst_idx = mad24(y, dst_step, dst_offset + x);
-
-#ifdef REVERSE
- dst[dst_idx] = src[src_idx + 2];
- dst[dst_idx + 1] = src[src_idx + 1];
- dst[dst_idx + 2] = src[src_idx];
-#elif defined ORDER
- dst[dst_idx] = src[src_idx];
- dst[dst_idx + 1] = src[src_idx + 1];
- dst[dst_idx + 2] = src[src_idx + 2];
-#endif
-
-#if dcn == 4
-#if scn == 3
- dst[dst_idx + 3] = MAX_NUM;
-#else
- dst[dst_idx + 3] = src[src_idx + 3];
-#endif
-#endif
- }
-}
///////////////////////////////////// RGB5x5 <-> RGB //////////////////////////////////////