diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index b25b71feee..aed9a7f0cf 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -50,7 +50,7 @@ using namespace cv; using namespace cv::ocl; -static void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx) +static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName) { int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); @@ -58,8 +58,8 @@ static void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx) std::string build_options = format("-D DEPTH_%d", src.depth()); vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); @@ -68,11 +68,11 @@ static void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx) args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - size_t gt[3] = { src.cols, src.rows, 1 }, lt[3] = { 16, 16, 1 }; - openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2Gray", gt, lt, args, -1, -1, build_options.c_str()); + size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; + openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } -static void Gray2RGB_caller(const oclMat &src, oclMat &dst) +static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName) { int channels = dst.channels(); std::string build_options = format("-D DEPTH_%d", src.depth()); @@ -80,99 +80,8 @@ static void Gray2RGB_caller(const oclMat &src, oclMat &dst) int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "Gray2RGB", gt, lt, args, -1, -1, build_options.c_str()); -} - -static void RGB2YUV_caller(const oclMat &src, oclMat &dst, int bidx) -{ - std::string build_options = format("-D DEPTH_%d", src.depth()); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2YUV", gt, lt, args, -1, -1, build_options.c_str()); -} - -static void YUV2RGB_caller(const oclMat &src, oclMat &dst, int bidx) -{ - int channels = dst.channels(); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - std::string buildOptions = format("-D DEPTH_%d", src.depth()); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGB", gt, lt, args, -1, -1, buildOptions.c_str()); -} - -static void YUV2RGB_NV12_caller(const oclMat &src, oclMat &dst, int bidx) -{ - std::string build_options = format("-D DEPTH_%d", src.depth()); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = {dst.cols / 2, dst.rows / 2, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGBA_NV12", gt, lt, args, -1, -1, build_options.c_str()); -} - -static void YCrCb2RGB_caller(const oclMat &src, oclMat &dst, int bidx) -{ - int channels = dst.channels(); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - std::string buildOptions = format("-D DEPTH_%d", src.depth()); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); @@ -182,29 +91,8 @@ static void YCrCb2RGB_caller(const oclMat &src, oclMat &dst, int bidx) args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - size_t gt[3] = { src.cols, src.rows, 1 }, lt[3] = { 16, 16, 1 }; - openCLExecuteKernel(src.clCxt, &cvt_color, "YCrCb2RGB", gt, lt, args, -1, -1, buildOptions.c_str()); -} - -static void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx) -{ - std::string build_options = format("-D DEPTH_%d", src.depth()); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2YCrCb", gt, lt, args, -1, -1, build_options.c_str()); + size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; + openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) @@ -232,7 +120,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, 1)); - RGB2Gray_caller(src, dst, bidx); + fromRGB_caller(src, dst, bidx, "RGB2Gray"); break; } case CV_GRAY2BGR: @@ -241,7 +129,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 1); dcn = code == CV_GRAY2BGRA ? 4 : 3; dst.create(sz, CV_MAKETYPE(depth, dcn)); - Gray2RGB_caller(src, dst); + toRGB_caller(src, dst, 0, "Gray2RGB"); break; } case CV_BGR2YUV: @@ -250,7 +138,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2YUV ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, 3)); - RGB2YUV_caller(src, dst, bidx); + fromRGB_caller(src, dst, bidx, "RGB2YUV"); break; } case CV_YUV2BGR: @@ -261,7 +149,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); bidx = code == CV_YUV2BGR ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, dcn)); - YUV2RGB_caller(src, dst, bidx); + toRGB_caller(src, dst, bidx, "YUV2RGB"); break; } case CV_YUV2RGB_NV12: @@ -276,7 +164,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) Size dstSz(sz.width, sz.height * 2 / 3); dst.create(dstSz, CV_MAKETYPE(depth, dcn)); - YUV2RGB_NV12_caller(src, dst, bidx); + toRGB_caller(src, dst, bidx, "YUV2RGBA_NV12"); break; } case CV_BGR2YCrCb: @@ -285,7 +173,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2YCrCb ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, 3)); - RGB2YCrCb_caller(src, dst, bidx); + fromRGB_caller(src, dst, bidx, "RGB2YCrCb"); break; } case CV_YCrCb2BGR: @@ -296,7 +184,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); bidx = code == CV_YCrCb2BGR ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, dcn)); - YCrCb2RGB_caller(src, dst, bidx); + toRGB_caller(src, dst, bidx, "YCrCb2RGB"); break; } /* diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index de53da52ee..d9426b28b1 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -100,7 +100,7 @@ __kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step, } } -__kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int channels, +__kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int channels, int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst, int src_offset, int dst_offset) { @@ -203,17 +203,17 @@ __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, int width, int height, __global const uchar* src, __global uchar* dst, +__kernel void YUV2RGBA_NV12(int cols, int rows, int src_step, int dst_step, int channels, + int bidx, __global const uchar* src, __global uchar* dst, int src_offset, int dst_offset) { - const int x = get_global_id(0); // max_x = width / 2 - const int y = get_global_id(1); // max_y = height/ 2 + const int x = get_global_id(0); + const int y = get_global_id(1); - if (y < height / 2 && x < width / 2 ) + 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(height + y, 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); diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index b2f5c6fb2d..df52b94b27 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -116,6 +116,8 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) #define CVTCODE(name) cv::COLOR_ ## name +// RGB <-> Gray + OCL_TEST_P(CvtColor, RGB2GRAY) { doTest(3, 1, CVTCODE(RGB2GRAY)); @@ -152,6 +154,7 @@ OCL_TEST_P(CvtColor, GRAY2BGRA) doTest(1, 4, CVTCODE(GRAY2BGRA)); } +// RGB <-> YUV OCL_TEST_P(CvtColor, RGB2YUV) { @@ -186,6 +189,7 @@ OCL_TEST_P(CvtColor, YUV2BGRA) doTest(3, 4, CVTCODE(YUV2BGR)); } +// RGB <-> YCrCb OCL_TEST_P(CvtColor, RGB2YCrCb) { @@ -220,6 +224,8 @@ OCL_TEST_P(CvtColor, YCrCb2BGRA) doTest(3, 4, CVTCODE(YCrCb2BGR)); } +// YUV -> RGBA_NV12 + struct CvtColor_YUV420 : public CvtColor { @@ -262,7 +268,6 @@ OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) doTest(1, 3, CV_YUV2BGR_NV12); } - INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor, testing::Combine( testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)),