From c8707b891b1cd7ca9f717ce95598cdab26865aa3 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Mon, 6 Oct 2014 19:19:44 +0400 Subject: [PATCH] Added OCL code for RGB[A]|BGR[A] -> YUV_[YV12|IYUV] color conversion --- modules/imgproc/src/color.cpp | 19 +++++++ modules/imgproc/src/opencl/cvtcolor.cl | 73 +++++++++++++++++++++++- modules/imgproc/test/ocl/test_color.cpp | 76 ++++++++++++++++++------- 3 files changed, 148 insertions(+), 20 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index a6cebe6bc0..6cca73652f 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -5011,6 +5011,25 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) src.rowRange(0, dstSz.height).copyTo(dst); return true; } + case COLOR_RGB2YUV_YV12: case COLOR_BGR2YUV_YV12: case COLOR_RGBA2YUV_YV12: case COLOR_BGRA2YUV_YV12: + case COLOR_RGB2YUV_IYUV: case COLOR_BGR2YUV_IYUV: case COLOR_RGBA2YUV_IYUV: case COLOR_BGRA2YUV_IYUV: + { + if (dcn <= 0) dcn = 1; + bidx = code == COLOR_BGRA2YUV_YV12 || code == COLOR_BGR2YUV_YV12 || + code == COLOR_BGRA2YUV_IYUV || code == COLOR_BGR2YUV_IYUV ? 0 : 2; + uidx = code == COLOR_RGBA2YUV_YV12 || code == COLOR_RGB2YUV_YV12 || + code == COLOR_BGRA2YUV_YV12 || code == COLOR_BGR2YUV_YV12 ? 1 : 0; + + CV_Assert( (scn == 3 || scn == 4) && depth == CV_8U ); + CV_Assert( dcn == 1 ); + CV_Assert( sz.width % 2 == 0 && sz.height % 2 == 0 ); + + dstSz = Size(sz.width, sz.height / 2 * 3); + globalsize[0] = dstSz.width / 2; globalsize[1] = (dstSz.height/3 + pxPerWIy - 1) / pxPerWIy; + k.create("RGB2YUV_YV12_IYUV", ocl::imgproc::cvtcolor_oclsrc, + opts + format("-D dcn=%d -D bidx=%d -D uidx=%d", dcn, bidx, uidx)); + break; + } case COLOR_BGR2YCrCb: case COLOR_RGB2YCrCb: { diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 4ac516ae8c..f57c0c068c 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -115,6 +115,10 @@ enum #define uidx 0 #endif +#ifndef yidx +#define yidx 0 +#endif + #define __CAT(x, y) x##y #define CAT(x, y) __CAT(x, y) @@ -317,7 +321,7 @@ __kernel void YUV2RGB_NV(__global const uchar* srcptr, int src_step, int src_off { __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]; @@ -446,6 +450,73 @@ __kernel void YUV2RGB_YV12_IYUV(__global const uchar* srcptr, int src_step, int } } +__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}; diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index 43343469c5..1484db9720 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -322,7 +322,7 @@ OCL_TEST_P(CvtColor8u32f, Luv2LRGBA) { performTest(3, 4, CVTCODE(Luv2LRGB), dept // YUV420 -> RGBA -struct CvtColor_YUV420 : +struct CvtColor_YUV2RGB_420 : public CvtColor { void generateTestData(int channelsIn, int channelsOut) @@ -344,24 +344,57 @@ struct CvtColor_YUV420 : } }; -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())); @@ -374,7 +407,12 @@ OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor, 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()));