From 8a236468973be7817231c2da73f61e0d27da0eaa Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Wed, 27 Nov 2013 20:00:35 +0400
Subject: [PATCH] RGB[A] <- RGB5x5
---
modules/imgproc/src/color.cpp | 16 ++++++-
modules/imgproc/src/opencl/cvtcolor.cl | 40 ++++++++--------
modules/imgproc/test/ocl/test_color.cpp | 26 +++++------
modules/ocl/src/opencl/cvt_color.cl | 61 +------------------------
4 files changed, 49 insertions(+), 94 deletions(-)
diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp
index c375b0533..88655757c 100644
--- a/modules/imgproc/src/color.cpp
+++ b/modules/imgproc/src/color.cpp
@@ -92,6 +92,7 @@
#include "precomp.hpp"
#include "opencl_kernels.hpp"
#include <limits>
+#include <iostream>
#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n))
@@ -2728,7 +2729,20 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
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));
+ format("-D depth=%d -D scn=2 -D dcn=%d -D bidx=%d -D greenbits=%d", depth, dcn, bidx, greenbits));
+ break;
+ }
+ case COLOR_BGR2BGR565: case COLOR_BGR2BGR555: case COLOR_RGB2BGR565: case COLOR_RGB2BGR555:
+ case COLOR_BGRA2BGR565: case COLOR_BGRA2BGR555: case COLOR_RGBA2BGR565: case COLOR_RGBA2BGR555:
+ {
+ CV_Assert((scn == 3 || scn == 4) && depth == CV_8U );
+ bidx = code == COLOR_BGR2BGR565 || code == COLOR_BGR2BGR555 ||
+ code == COLOR_BGRA2BGR565 || code == COLOR_BGRA2BGR555 ? 0 : 2;
+ int greenbits = code == COLOR_BGR2BGR565 || code == COLOR_RGB2BGR565 ||
+ code == COLOR_BGRA2BGR565 || code == COLOR_RGBA2BGR565 ? 6 : 5;
+ dcn = 2;
+ k.create("RGB2RGB5x5", ocl::imgproc::cvtcolor_oclsrc,
+ format("-D depth=%d -D scn=%d -D dcn=2 -D bidx=%d -D greenbits=%d", depth, scn, bidx, greenbits));
break;
}
case COLOR_BGR2GRAY: case COLOR_BGRA2GRAY:
diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl
index e6ceada8e..0d8272fae 100644
--- a/modules/imgproc/src/opencl/cvtcolor.cl
+++ b/modules/imgproc/src/opencl/cvtcolor.cl
@@ -490,28 +490,28 @@ __kernel void RGB5x52RGB(__global const uchar* src, int src_step, int src_offset
}
}
-//__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);
+__kernel void RGB2RGB5x5(__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 << 2));
-// int dst_idx = mad24(y, dst_step, dst_offset + x);
+ 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);
-//#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
-// }
-//}
+#if greenbits == 6
+ *((__global ushort*)(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
+ *((__global ushort*)(dst + dst_idx)) = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)|((src[src_idx + (bidx^2)]&~7) << 7));
+#else
+ *((__global ushort*)(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
+ }
+}
diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp
index d65c6f240..695d3d164 100644
--- a/modules/imgproc/test/ocl/test_color.cpp
+++ b/modules/imgproc/test/ocl/test_color.cpp
@@ -169,7 +169,7 @@ OCL_TEST_P(CvtColor, XYZ2BGRA) { performTest(3, 4, CVTCODE(XYZ2BGR)); }
// RGB <-> HSV
-//typedef CvtColor CvtColor8u32f;
+typedef CvtColor CvtColor8u32f;
//OCL_TEST_P(CvtColor8u32f, RGB2HSV) { performTest(3, 3, CVTCODE(RGB2HSV)); }
//OCL_TEST_P(CvtColor8u32f, BGR2HSV) { performTest(3, 3, CVTCODE(BGR2HSV)); }
@@ -227,15 +227,15 @@ 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)); }
-//OCL_TEST_P(CvtColor8u, BGRA2BGR565) { performTest(4, 2, CVTCODE(BGRA2BGR565)); }
-//OCL_TEST_P(CvtColor8u, RGBA2BGR565) { performTest(4, 2, CVTCODE(RGBA2BGR565)); }
+OCL_TEST_P(CvtColor8u, BGR2BGR565) { performTest(3, 2, CVTCODE(BGR2BGR565)); }
+OCL_TEST_P(CvtColor8u, RGB2BGR565) { performTest(3, 2, CVTCODE(RGB2BGR565)); }
+OCL_TEST_P(CvtColor8u, BGRA2BGR565) { performTest(4, 2, CVTCODE(BGRA2BGR565)); }
+OCL_TEST_P(CvtColor8u, RGBA2BGR565) { performTest(4, 2, CVTCODE(RGBA2BGR565)); }
-//OCL_TEST_P(CvtColor8u, BGR2BGR555) { performTest(3, 2, CVTCODE(BGR2BGR555)); }
-//OCL_TEST_P(CvtColor8u, RGB2BGR555) { performTest(3, 2, CVTCODE(RGB2BGR555)); }
-//OCL_TEST_P(CvtColor8u, BGRA2BGR555) { performTest(4, 2, CVTCODE(BGRA2BGR555)); }
-//OCL_TEST_P(CvtColor8u, RGBA2BGR555) { performTest(4, 2, CVTCODE(RGBA2BGR555)); }
+OCL_TEST_P(CvtColor8u, BGR2BGR555) { performTest(3, 2, CVTCODE(BGR2BGR555)); }
+OCL_TEST_P(CvtColor8u, RGB2BGR555) { performTest(3, 2, CVTCODE(RGB2BGR555)); }
+OCL_TEST_P(CvtColor8u, BGRA2BGR555) { performTest(4, 2, CVTCODE(BGRA2BGR555)); }
+OCL_TEST_P(CvtColor8u, RGBA2BGR555) { performTest(4, 2, CVTCODE(RGBA2BGR555)); }
// RGB5x5 <-> Gray
@@ -280,11 +280,11 @@ OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { performTest(1, 3, COLOR_YUV2RGB_NV12
OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { performTest(1, 3, COLOR_YUV2BGR_NV12); }
-//OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u,
-// testing::Combine(testing::Values(MatDepth(CV_8U)), Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u,
+ testing::Combine(testing::Values(MatDepth(CV_8U)), Bool()));
-//OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u32f,
-// testing::Combine(testing::Values(MatDepth(CV_8U), MatDepth(CV_32F)), Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u32f,
+ testing::Combine(testing::Values(MatDepth(CV_8U), MatDepth(CV_32F)), Bool()));
OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor,
testing::Combine(
diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl
index 5786505bd..3b2e84c0a 100644
--- a/modules/ocl/src/opencl/cvt_color.cl
+++ b/modules/ocl/src/opencl/cvt_color.cl
@@ -91,66 +91,7 @@ enum
BLOCK_SIZE = 256
};
-
-///////////////////////////////////// RGB5x5 <-> RGB //////////////////////////////////////
-
-__kernel void RGB5x52RGB(int cols, int rows, int src_step, int dst_step, int bidx,
- __global const ushort * src, __global uchar * 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);
- int dst_idx = mad24(y, dst_step, dst_offset + (x << 2));
- ushort t = 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 //////////////////////////////////////
+///////////////////////////////////// RGB5x5 <-> Gray //////////////////////////////////////
__kernel void BGR5x52Gray(int cols, int rows, int src_step, int dst_step, int bidx,
__global const ushort * src, __global uchar * dst,