#define uidx 0
#endif
+#ifndef yidx
+#define yidx 0
+#endif
+
#define __CAT(x, y) x##y
#define CAT(x, y) __CAT(x, y)
{
__global const uchar* ysrc = srcptr + mad24(y << 1, src_step, (x << 1) + src_offset);
__global const uchar* usrc = srcptr + mad24(rows + y, src_step, (x << 1) + src_offset);
- __global uchar* dst1 = dstptr + mad24(y << 1, dst_step, x * (dcn<<1) + dt_offset);
+ __global uchar* dst1 = dstptr + mad24(y << 1, dst_step, mad24(x, dcn<<1, dt_offset));
__global uchar* dst2 = dst1 + dst_step;
int Y1 = ysrc[0];
}
}
+__constant int ITUR_BT_601_CRY = 269484;
+__constant int ITUR_BT_601_CGY = 528482;
+__constant int ITUR_BT_601_CBY = 102760;
+__constant int ITUR_BT_601_CRU = -155188;
+__constant int ITUR_BT_601_CGU = -305135;
+__constant int ITUR_BT_601_CBU = 460324;
+__constant int ITUR_BT_601_CGV = -385875;
+__constant int ITUR_BT_601_CBV = -74448;
+__constant int YSHIFT = 17301504;
+__constant int UVSHIFT = 134742016;
+
+__kernel void RGB2YUV_YV12_IYUV(__global const uchar* srcptr, int src_step, int src_offset,
+ __global uchar* dstptr, int dst_step, int dst_offset,
+ int rows, int cols)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1) * PIX_PER_WI_Y;
+
+ if (x < cols/2)
+ {
+ int src_index = mad24(y << 1, src_step, mad24(x << 1, scn, src_offset));
+ int ydst_index = mad24(y << 1, dst_step, (x << 1) + dst_offset);
+ int y_rows = rows / 3 * 2;
+ int vsteps[2] = { cols >> 1, dst_step - (cols >> 1)};
+
+ #pragma unroll
+ for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
+ {
+ if (y < rows / 3)
+ {
+ __global const uchar* src1 = srcptr + src_index;
+ __global const uchar* src2 = src1 + src_step;
+ __global uchar* ydst1 = dstptr + ydst_index;
+ __global uchar* ydst2 = ydst1 + dst_step;
+
+ __global uchar* udst = dstptr + mad24(y_rows + (y>>1), dst_step, dst_offset + (y%2)*(cols >> 1) + x);
+ __global uchar* vdst = udst + mad24(y_rows >> 2, dst_step, y_rows % 4 ? vsteps[y%2] : 0);
+
+ int4 src_pix1 = convert_int4(vload4(0, src1));
+ int4 src_pix2 = convert_int4(vload4(0, src1+scn));
+ int4 src_pix3 = convert_int4(vload4(0, src2));
+ int4 src_pix4 = convert_int4(vload4(0, src2+scn));
+
+ int y00 = mad24(ITUR_BT_601_CRY, src_pix1.R_COMP, mad24(ITUR_BT_601_CGY, src_pix1.G_COMP, mad24(ITUR_BT_601_CBY, src_pix1.B_COMP, YSHIFT)));
+ int y01 = mad24(ITUR_BT_601_CRY, src_pix2.R_COMP, mad24(ITUR_BT_601_CGY, src_pix2.G_COMP, mad24(ITUR_BT_601_CBY, src_pix2.B_COMP, YSHIFT)));
+ int y10 = mad24(ITUR_BT_601_CRY, src_pix3.R_COMP, mad24(ITUR_BT_601_CGY, src_pix3.G_COMP, mad24(ITUR_BT_601_CBY, src_pix3.B_COMP, YSHIFT)));
+ int y11 = mad24(ITUR_BT_601_CRY, src_pix4.R_COMP, mad24(ITUR_BT_601_CGY, src_pix4.G_COMP, mad24(ITUR_BT_601_CBY, src_pix4.B_COMP, YSHIFT)));
+
+ ydst1[0] = convert_uchar_sat(y00 >> ITUR_BT_601_SHIFT);
+ ydst1[1] = convert_uchar_sat(y01 >> ITUR_BT_601_SHIFT);
+ ydst2[0] = convert_uchar_sat(y10 >> ITUR_BT_601_SHIFT);
+ ydst2[1] = convert_uchar_sat(y11 >> ITUR_BT_601_SHIFT);
+
+ int uv[2] = { mad24(ITUR_BT_601_CRU, src_pix1.R_COMP, mad24(ITUR_BT_601_CGU, src_pix1.G_COMP, mad24(ITUR_BT_601_CBU, src_pix1.B_COMP, UVSHIFT))),
+ mad24(ITUR_BT_601_CBU, src_pix1.R_COMP, mad24(ITUR_BT_601_CGV, src_pix1.G_COMP, mad24(ITUR_BT_601_CBV, src_pix1.B_COMP, UVSHIFT))) };
+
+ udst[0] = convert_uchar_sat(uv[uidx] >> ITUR_BT_601_SHIFT);
+ vdst[0] = convert_uchar_sat(uv[1-uidx] >> ITUR_BT_601_SHIFT);
+
+ ++y;
+ src_index += 2*src_step;
+ ydst_index += 2*dst_step;
+ }
+ }
+ }
+}
+
///////////////////////////////////// RGB <-> YCrCb //////////////////////////////////////
__constant float c_RGB2YCrCbCoeffs_f[5] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};
// YUV420 -> RGBA
-struct CvtColor_YUV420 :
+struct CvtColor_YUV2RGB_420 :
public CvtColor
{
void generateTestData(int channelsIn, int channelsOut)
}
};
-OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) { performTest(1, 4, CVTCODE(YUV2RGBA_NV12)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) { performTest(1, 4, CVTCODE(YUV2BGRA_NV12)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { performTest(1, 3, CVTCODE(YUV2RGB_NV12)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { performTest(1, 3, CVTCODE(YUV2BGR_NV12)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV21) { performTest(1, 4, CVTCODE(YUV2RGBA_NV21)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV21) { performTest(1, 4, CVTCODE(YUV2BGRA_NV21)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV21) { performTest(1, 3, CVTCODE(YUV2RGB_NV21)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV21) { performTest(1, 3, CVTCODE(YUV2BGR_NV21)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_YV12) { performTest(1, 4, CVTCODE(YUV2RGBA_YV12)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_YV12) { performTest(1, 4, CVTCODE(YUV2BGRA_YV12)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2RGB_YV12) { performTest(1, 3, CVTCODE(YUV2RGB_YV12)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2BGR_YV12) { performTest(1, 3, CVTCODE(YUV2BGR_YV12)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_IYUV) { performTest(1, 4, CVTCODE(YUV2RGBA_IYUV)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_IYUV) { performTest(1, 4, CVTCODE(YUV2BGRA_IYUV)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2RGB_IYUV) { performTest(1, 3, CVTCODE(YUV2RGB_IYUV)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2BGR_IYUV) { performTest(1, 3, CVTCODE(YUV2BGR_IYUV)); }
-OCL_TEST_P(CvtColor_YUV420, YUV2GRAY_420) { performTest(1, 1, CVTCODE(YUV2GRAY_420)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2RGBA_NV12) { performTest(1, 4, CVTCODE(YUV2RGBA_NV12)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2BGRA_NV12) { performTest(1, 4, CVTCODE(YUV2BGRA_NV12)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2RGB_NV12) { performTest(1, 3, CVTCODE(YUV2RGB_NV12)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2BGR_NV12) { performTest(1, 3, CVTCODE(YUV2BGR_NV12)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2RGBA_NV21) { performTest(1, 4, CVTCODE(YUV2RGBA_NV21)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2BGRA_NV21) { performTest(1, 4, CVTCODE(YUV2BGRA_NV21)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2RGB_NV21) { performTest(1, 3, CVTCODE(YUV2RGB_NV21)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2BGR_NV21) { performTest(1, 3, CVTCODE(YUV2BGR_NV21)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2RGBA_YV12) { performTest(1, 4, CVTCODE(YUV2RGBA_YV12)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2BGRA_YV12) { performTest(1, 4, CVTCODE(YUV2BGRA_YV12)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2RGB_YV12) { performTest(1, 3, CVTCODE(YUV2RGB_YV12)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2BGR_YV12) { performTest(1, 3, CVTCODE(YUV2BGR_YV12)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2RGBA_IYUV) { performTest(1, 4, CVTCODE(YUV2RGBA_IYUV)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2BGRA_IYUV) { performTest(1, 4, CVTCODE(YUV2BGRA_IYUV)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2RGB_IYUV) { performTest(1, 3, CVTCODE(YUV2RGB_IYUV)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2BGR_IYUV) { performTest(1, 3, CVTCODE(YUV2BGR_IYUV)); }
+OCL_TEST_P(CvtColor_YUV2RGB_420, YUV2GRAY_420) { performTest(1, 1, CVTCODE(YUV2GRAY_420)); }
+
+// RGBA -> YUV420
+
+struct CvtColor_RGB2YUV_420 :
+ public CvtColor
+{
+ void generateTestData(int channelsIn, int channelsOut)
+ {
+ const int srcType = CV_MAKE_TYPE(depth, channelsIn);
+ const int dstType = CV_MAKE_TYPE(depth, channelsOut);
+ Size srcRoiSize = randomSize(1, MAX_VALUE);
+ srcRoiSize.width *= 2;
+ srcRoiSize.height *= 2;
+ Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(src, src_roi, srcRoiSize, srcBorder, srcType, 2, 100);
+
+ Size dstRoiSize(srcRoiSize.width, srcRoiSize.height / 2 * 3);
+ Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(dst, dst_roi, dstRoiSize, dstBorder, dstType, 5, 16);
+
+ UMAT_UPLOAD_INPUT_PARAMETER(src);
+ UMAT_UPLOAD_OUTPUT_PARAMETER(dst);
+ }
+};
+
+OCL_TEST_P(CvtColor_RGB2YUV_420, RGBA2YUV_YV12) { performTest(4, 1, CVTCODE(RGBA2YUV_YV12)); }
+OCL_TEST_P(CvtColor_RGB2YUV_420, BGRA2YUV_YV12) { performTest(4, 1, CVTCODE(BGRA2YUV_YV12)); }
+OCL_TEST_P(CvtColor_RGB2YUV_420, RGB2YUV_YV12) { performTest(3, 1, CVTCODE(RGB2YUV_YV12)); }
+OCL_TEST_P(CvtColor_RGB2YUV_420, BGR2YUV_YV12) { performTest(3, 1, CVTCODE(BGR2YUV_YV12)); }
+OCL_TEST_P(CvtColor_RGB2YUV_420, RGBA2YUV_IYUV) { performTest(4, 1, CVTCODE(RGBA2YUV_IYUV)); }
+OCL_TEST_P(CvtColor_RGB2YUV_420, BGRA2YUV_IYUV) { performTest(4, 1, CVTCODE(BGRA2YUV_IYUV)); }
+OCL_TEST_P(CvtColor_RGB2YUV_420, RGB2YUV_IYUV) { performTest(3, 1, CVTCODE(RGB2YUV_IYUV)); }
+OCL_TEST_P(CvtColor_RGB2YUV_420, BGR2YUV_IYUV) { performTest(3, 1, CVTCODE(BGR2YUV_IYUV)); }
OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u,
testing::Combine(testing::Values(MatDepth(CV_8U)), Bool()));
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)),
Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor_YUV420,
+OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor_YUV2RGB_420,
+ testing::Combine(
+ testing::Values(MatDepth(CV_8U)),
+ Bool()));
+
+OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor_RGB2YUV_420,
testing::Combine(
testing::Values(MatDepth(CV_8U)),
Bool()));