mirror of
https://github.com/opencv/opencv.git
synced 2024-11-29 13:47:32 +08:00
RGB[A] <- RGB5x5
This commit is contained in:
parent
81b9c9c104
commit
8a23646897
@ -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:
|
||||
|
@ -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
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
@ -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(
|
||||
|
@ -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,
|
||||
|
Loading…
Reference in New Issue
Block a user