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_BGR2GRAY:
- case COLOR_BGRA2GRAY:
- case COLOR_RGB2GRAY:
- case COLOR_RGBA2GRAY:
+ case COLOR_BGR2GRAY: case COLOR_BGRA2GRAY:
+ case COLOR_RGB2GRAY: case COLOR_RGBA2GRAY:
{
CV_Assert(scn == 3 || scn == 4);
bidx = code == COLOR_BGR2GRAY || code == COLOR_BGRA2GRAY ? 0 : 2;
format("-D depth=%d -D scn=3 -D dcn=%d -D bidx=%d", depth, dcn, bidx));
break;
}
- case COLOR_YUV2RGB_NV12:
- case COLOR_YUV2BGR_NV12:
- case COLOR_YUV2RGBA_NV12:
- case COLOR_YUV2BGRA_NV12:
+ case COLOR_YUV2RGB_NV12: case COLOR_YUV2BGR_NV12:
+ case COLOR_YUV2RGBA_NV12: case COLOR_YUV2BGRA_NV12:
{
CV_Assert( scn == 1 );
CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U );
}
case COLOR_YCrCb2BGR:
case COLOR_YCrCb2RGB:
+ {
+ if( dcn <= 0 )
+ dcn = 3;
+ CV_Assert(scn == 3 && (dcn == 3 || dcn == 4));
+ bidx = code == COLOR_YCrCb2BGR ? 0 : 2;
+ k.create("YCrCb2RGB", ocl::imgproc::cvtcolor_oclsrc,
+ format("-D depth=%d -D scn=%d -D dcn=%d -D bidx=%d", depth, scn, dcn, bidx));
break;
+ }
/*
case COLOR_BGR5652GRAY: case COLOR_BGR5552GRAY:
case COLOR_GRAY2BGR565: case COLOR_GRAY2BGR555:
- case COLOR_BGR2YCrCb: case COLOR_RGB2YCrCb:
case COLOR_BGR2XYZ: case COLOR_RGB2XYZ:
case COLOR_XYZ2BGR: case COLOR_XYZ2RGB:
case COLOR_BGR2HSV: case COLOR_RGB2HSV: case COLOR_BGR2HSV_FULL: case COLOR_RGB2HSV_FULL:
int stype = _src.type();
int scn = CV_MAT_CN(stype), depth = CV_MAT_DEPTH(stype), bidx;
- if( use_opencl && ocl_cvtColor(_src, _dst, code, dcn) )
- return;
+ if( use_opencl /*&& ocl_cvtColor(_src, _dst, code, dcn)*/ )
+ return (void)ocl_cvtColor(_src, _dst, code, dcn);
Mat src = _src.getMat(), dst;
Size sz = src.size();
}
}
-///////////////////////////////////// RGB -> YCrCb //////////////////////////////////////
+///////////////////////////////////// RGB <-> YCrCb //////////////////////////////////////
__constant float c_RGB2YCrCbCoeffs_f[5] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};
__constant int c_RGB2YCrCbCoeffs_i[5] = {R2Y, G2Y, B2Y, 11682, 9241};
dst[2] = SAT_CAST( Cb );
}
}
+
+__constant float c_YCrCb2RGBCoeffs_f[4] = { 1.403f, -0.714f, -0.344f, 1.773f };
+__constant int c_YCrCb2RGBCoeffs_i[4] = { 22987, -11698, -5636, 29049 };
+
+__kernel void YCrCb2RGB(__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);
+ __global const DATA_TYPE * srcptr = (__global const DATA_TYPE*)(src + src_idx);
+ __global DATA_TYPE * dstptr = (__global DATA_TYPE*)(dst + dst_idx);
+
+ DATA_TYPE y = srcptr[0], cr = srcptr[1], cb = srcptr[2];
+
+#ifdef DEPTH_5
+ __constant float * coeff = c_YCrCb2RGBCoeffs_f;
+ float r = y + coeff[0] * (cr - HALF_MAX);
+ float g = y + coeff[1] * (cr - HALF_MAX) + coeff[2] * (cb - HALF_MAX);
+ float b = y + coeff[3] * (cb - HALF_MAX);
+#else
+ __constant int * coeff = c_YCrCb2RGBCoeffs_i;
+ int r = y + CV_DESCALE(coeff[0] * (cr - HALF_MAX), yuv_shift);
+ int g = y + CV_DESCALE(coeff[1] * (cr - HALF_MAX) + coeff[2] * (cb - HALF_MAX), yuv_shift);
+ int b = y + CV_DESCALE(coeff[3] * (cb - HALF_MAX), yuv_shift);
+#endif
+
+ dstptr[(bidx^2)] = SAT_CAST(r);
+ dstptr[1] = SAT_CAST(g);
+ dstptr[bidx] = SAT_CAST(b);
+#if dcn == 4
+ dstptr[3] = MAX_NUM;
+#endif
+ }
+}
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
BLOCK_SIZE = 256
};
-///////////////////////////////////// RGB <-> GRAY //////////////////////////////////////
-
-__kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step,
- int bidx, __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)
- {
- int src_idx = mad24(y, src_step, src_offset + (x << 2));
- int dst_idx = mad24(y, dst_step, dst_offset + x);
-#ifdef DEPTH_5
- dst[dst_idx] = src[src_idx + bidx] * 0.114f + src[src_idx + 1] * 0.587f + src[src_idx + (bidx^2)] * 0.299f;
-#else
- dst[dst_idx] = (DATA_TYPE)CV_DESCALE((src[src_idx + bidx] * B2Y + src[src_idx + 1] * G2Y + src[src_idx + (bidx^2)] * R2Y), yuv_shift);
-#endif
- }
-}
-
-__kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int bidx,
- __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)
- {
- int src_idx = mad24(y, src_step, src_offset + x);
- int dst_idx = mad24(y, dst_step, dst_offset + (x << 2));
-
- DATA_TYPE val = src[src_idx];
- dst[dst_idx] = val;
- dst[dst_idx + 1] = val;
- dst[dst_idx + 2] = val;
-#if dcn == 4
- dst[dst_idx + 3] = MAX_NUM;
-#endif
- }
-}
-
-///////////////////////////////////// RGB <-> YUV //////////////////////////////////////
-
-__constant float c_RGB2YUVCoeffs_f[5] = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f };
-__constant int c_RGB2YUVCoeffs_i[5] = { B2Y, G2Y, R2Y, 8061, 14369 };
-
-__kernel void RGB2YUV(int cols, int rows, int src_step, int dst_step,
- int bidx, __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);
- DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] };
-
-#ifdef DEPTH_5
- __constant float * coeffs = c_RGB2YUVCoeffs_f;
- DATA_TYPE Y = rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx];
- DATA_TYPE Cr = (rgb[bidx^2] - Y) * coeffs[3] + HALF_MAX;
- DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX;
-#else
- __constant int * coeffs = c_RGB2YUVCoeffs_i;
- int delta = HALF_MAX * (1 << yuv_shift);
- int Y = CV_DESCALE(rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx], yuv_shift);
- int Cr = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[3] + delta, yuv_shift);
- int Cb = CV_DESCALE((rgb[bidx] - Y) * coeffs[4] + delta, yuv_shift);
-#endif
-
- dst[dst_idx] = SAT_CAST( Y );
- dst[dst_idx + 1] = SAT_CAST( Cr );
- dst[dst_idx + 2] = SAT_CAST( Cb );
- }
-}
-
-__constant float c_YUV2RGBCoeffs_f[5] = { 2.032f, -0.395f, -0.581f, 1.140f };
-__constant int c_YUV2RGBCoeffs_i[5] = { 33292, -6472, -9519, 18678 };
-
-__kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step,
- int bidx, __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);
- DATA_TYPE yuv[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] };
-
-#ifdef DEPTH_5
- __constant float * coeffs = c_YUV2RGBCoeffs_f;
- float b = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[3];
- float g = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1];
- float r = yuv[0] + (yuv[1] - HALF_MAX) * coeffs[0];
-#else
- __constant int * coeffs = c_YUV2RGBCoeffs_i;
- int b = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[3], yuv_shift);
- int g = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1], yuv_shift);
- int r = yuv[0] + CV_DESCALE((yuv[1] - HALF_MAX) * coeffs[0], yuv_shift);
-#endif
-
- dst[dst_idx + bidx] = SAT_CAST( b );
- dst[dst_idx + 1] = SAT_CAST( g );
- dst[dst_idx + (bidx^2)] = SAT_CAST( r );
-#if dcn == 4
- dst[dst_idx + 3] = MAX_NUM;
-#endif
- }
-}
-
-__constant int ITUR_BT_601_CY = 1220542;
-__constant int ITUR_BT_601_CUB = 2116026;
-__constant int ITUR_BT_601_CUG = 409993;
-__constant int ITUR_BT_601_CVG = 852492;
-__constant int ITUR_BT_601_CVR = 1673527;
-__constant int ITUR_BT_601_SHIFT = 20;
-
-__kernel void YUV2RGBA_NV12(int cols, int rows, int src_step, int dst_step,
- int bidx, __global const uchar* src, __global uchar* dst,
- int src_offset, int dst_offset)
-{
- const int x = get_global_id(0);
- const int y = get_global_id(1);
-
- if (y < rows / 2 && x < cols / 2 )
- {
- __global const uchar* ysrc = src + mad24(y << 1, src_step, (x << 1) + src_offset);
- __global const uchar* usrc = src + mad24(rows + y, src_step, (x << 1) + src_offset);
- __global uchar* dst1 = dst + mad24(y << 1, dst_step, (x << 3) + dst_offset);
- __global uchar* dst2 = dst + mad24((y << 1) + 1, dst_step, (x << 3) + dst_offset);
-
- int Y1 = ysrc[0];
- int Y2 = ysrc[1];
- int Y3 = ysrc[src_step];
- int Y4 = ysrc[src_step + 1];
-
- int U = usrc[0] - 128;
- int V = usrc[1] - 128;
-
- int ruv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CVR * V;
- int guv = (1 << (ITUR_BT_601_SHIFT - 1)) - ITUR_BT_601_CVG * V - ITUR_BT_601_CUG * U;
- int buv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CUB * U;
-
- Y1 = max(0, Y1 - 16) * ITUR_BT_601_CY;
- dst1[2 - bidx] = convert_uchar_sat((Y1 + ruv) >> ITUR_BT_601_SHIFT);
- dst1[1] = convert_uchar_sat((Y1 + guv) >> ITUR_BT_601_SHIFT);
- dst1[bidx] = convert_uchar_sat((Y1 + buv) >> ITUR_BT_601_SHIFT);
- dst1[3] = 255;
-
- Y2 = max(0, Y2 - 16) * ITUR_BT_601_CY;
- dst1[6 - bidx] = convert_uchar_sat((Y2 + ruv) >> ITUR_BT_601_SHIFT);
- dst1[5] = convert_uchar_sat((Y2 + guv) >> ITUR_BT_601_SHIFT);
- dst1[4 + bidx] = convert_uchar_sat((Y2 + buv) >> ITUR_BT_601_SHIFT);
- dst1[7] = 255;
-
- Y3 = max(0, Y3 - 16) * ITUR_BT_601_CY;
- dst2[2 - bidx] = convert_uchar_sat((Y3 + ruv) >> ITUR_BT_601_SHIFT);
- dst2[1] = convert_uchar_sat((Y3 + guv) >> ITUR_BT_601_SHIFT);
- dst2[bidx] = convert_uchar_sat((Y3 + buv) >> ITUR_BT_601_SHIFT);
- dst2[3] = 255;
-
- Y4 = max(0, Y4 - 16) * ITUR_BT_601_CY;
- dst2[6 - bidx] = convert_uchar_sat((Y4 + ruv) >> ITUR_BT_601_SHIFT);
- dst2[5] = convert_uchar_sat((Y4 + guv) >> ITUR_BT_601_SHIFT);
- dst2[4 + bidx] = convert_uchar_sat((Y4 + buv) >> ITUR_BT_601_SHIFT);
- dst2[7] = 255;
- }
-}
-
-///////////////////////////////////// RGB <-> YCrCb //////////////////////////////////////
-
-__constant float c_RGB2YCrCbCoeffs_f[5] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};
-__constant int c_RGB2YCrCbCoeffs_i[5] = {R2Y, G2Y, B2Y, 11682, 9241};
-
-__kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step,
- int bidx, __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);
-
- DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] };
-
-#ifdef DEPTH_5
- __constant float * coeffs = c_RGB2YCrCbCoeffs_f;
- DATA_TYPE Y = rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx];
- DATA_TYPE Cr = (rgb[bidx^2] - Y) * coeffs[3] + HALF_MAX;
- DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX;
-#else
- __constant int * coeffs = c_RGB2YCrCbCoeffs_i;
- int delta = HALF_MAX * (1 << yuv_shift);
- int Y = CV_DESCALE(rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx], yuv_shift);
- int Cr = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[3] + delta, yuv_shift);
- int Cb = CV_DESCALE((rgb[bidx] - Y) * coeffs[4] + delta, yuv_shift);
-#endif
-
- dst[dst_idx] = SAT_CAST( Y );
- dst[dst_idx + 1] = SAT_CAST( Cr );
- dst[dst_idx + 2] = SAT_CAST( Cb );
- }
-}
__constant float c_YCrCb2RGBCoeffs_f[4] = { 1.403f, -0.714f, -0.344f, 1.773f };
__constant int c_YCrCb2RGBCoeffs_i[4] = { 22987, -11698, -5636, 29049 };