diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index 27c2fd5f05..b01dcef1f4 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -50,180 +50,189 @@ using namespace cv; using namespace cv::ocl; -#ifndef CV_DESCALE -#define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n)) -#endif - -#ifndef FLT_EPSILON -#define FLT_EPSILON 1.192092896e-07F -#endif - -namespace +static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName, + const oclMat & data = oclMat()) { - -void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx) -{ - int channels = src.oclchannels(); int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); 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 *)&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, "RGB2Gray", gt, lt, args, -1, -1, build_options.c_str()); -} - -void Gray2RGB_caller(const oclMat &src, oclMat &dst) -{ - 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_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()); -} - -void RGB2YUV_caller(const oclMat &src, oclMat &dst, int bidx) -{ - int channels = src.oclchannels(); - 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 *)&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, "RGB2YUV", gt, lt, args, -1, -1, build_options.c_str()); -} - -void YUV2RGB_caller(const oclMat &src, oclMat &dst, int bidx) -{ - int channels = src.oclchannels(); - 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()); -} - -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()); -} - -void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx) -{ - int channels = src.oclchannels(); - 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 *)&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, "RGB2YCrCb", gt, lt, args, -1, -1, build_options.c_str()); + if (!data.empty()) + args.push_back( make_pair( sizeof(cl_mem) , (void *)&data.data )); + + 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()); } -void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) +static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName, + const oclMat & data = oclMat()) +{ + std::string build_options = format("-D DEPTH_%d -D dcn=%d", src.depth(), dst.channels()); + 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 *)&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)); + 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 )); + + if (!data.empty()) + args.push_back( make_pair( sizeof(cl_mem) , (void *)&data.data )); + + 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 RGB_caller(const oclMat &src, oclMat &dst, bool reverse) +{ + std::string build_options = format("-D DEPTH_%d -D dcn=%d -D scn=%d -D %s", src.depth(), + dst.channels(), src.channels(), reverse ? "REVERSE" : "ORDER"); + 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 *)&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_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, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; + openCLExecuteKernel(src.clCxt, &cvt_color, "RGB", gt, lt, args, -1, -1, build_options.c_str()); +} + +static void fromRGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits, const std::string & kernelName) +{ + std::string build_options = format("-D DEPTH_%d -D greenbits=%d -D dcn=%d", + src.depth(), greenbits, dst.channels()); + int src_offset = src.offset >> 1, src_step = src.step >> 1; + int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step / dst.elemSize1(); + + vector > args; + 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)); + 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, 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 toRGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits, const std::string & kernelName) +{ + std::string build_options = format("-D DEPTH_%d -D greenbits=%d -D scn=%d", + src.depth(), greenbits, src.channels()); + int src_offset = (int)src.offset, src_step = (int)src.step; + int dst_offset = dst.offset >> 1, dst_step = dst.step >> 1; + + vector > args; + 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)); + 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, 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) { Size sz = src.size(); - int scn = src.oclchannels(), depth = src.depth(), bidx; + int scn = src.channels(), depth = src.depth(), bidx; CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32F); switch (code) { - /* - case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR: - case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA: - case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555: - case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555: - case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: - case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: - */ - case CV_BGR2GRAY: - case CV_BGRA2GRAY: - case CV_RGB2GRAY: - case CV_RGBA2GRAY: + case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR: + case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA: + { + CV_Assert(scn == 3 || scn == 4); + dcn = code == CV_BGR2BGRA || code == CV_RGB2BGRA || code == CV_BGRA2RGBA ? 4 : 3; + bool reverse = !(code == CV_BGR2BGRA || code == CV_BGRA2BGR); + dst.create(sz, CV_MAKE_TYPE(depth, dcn)); + RGB_caller(src, dst, reverse); + break; + } + case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555: + case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555: + { + CV_Assert((scn == 3 || scn == 4) && depth == CV_8U ); + bidx = code == CV_BGR2BGR565 || code == CV_BGR2BGR555 || + code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2; + int greenbits = code == CV_BGR2BGR565 || code == CV_RGB2BGR565 || + code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5; + dst.create(sz, CV_8UC2); + toRGB5x5_caller(src, dst, bidx, greenbits, "RGB2RGB5x5"); + break; + } + case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: + case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: + { + dcn = code == CV_BGR5652BGRA || code == CV_BGR5552BGRA || code == CV_BGR5652RGBA || code == CV_BGR5552RGBA ? 4 : 3; + CV_Assert((dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U); + bidx = code == CV_BGR5652BGR || code == CV_BGR5552BGR || + code == CV_BGR5652BGRA || code == CV_BGR5552BGRA ? 0 : 2; + int greenbits = code == CV_BGR5652BGR || code == CV_BGR5652RGB || + code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5; + dst.create(sz, CV_MAKETYPE(depth, dcn)); + fromRGB5x5_caller(src, dst, bidx, greenbits, "RGB5x52RGB"); + break; + } + case CV_BGR5652GRAY: case CV_BGR5552GRAY: + { + CV_Assert(scn == 2 && depth == CV_8U); + dst.create(sz, CV_8UC1); + int greenbits = code == CV_BGR5652GRAY ? 6 : 5; + fromRGB5x5_caller(src, dst, -1, greenbits, "BGR5x52Gray"); + break; + } + case CV_GRAY2BGR565: case CV_GRAY2BGR555: + { + CV_Assert(scn == 1 && depth == CV_8U); + dst.create(sz, CV_8UC2); + int greenbits = code == CV_GRAY2BGR565 ? 6 : 5; + toRGB5x5_caller(src, dst, -1, greenbits, "Gray2BGR5x5"); + break; + } + case CV_RGB2GRAY: case CV_BGR2GRAY: + case CV_RGBA2GRAY: case CV_BGRA2GRAY: { 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: @@ -232,40 +241,40 @@ 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: case CV_RGB2YUV: { CV_Assert(scn == 3 || scn == 4); - bidx = code == CV_RGB2YUV ? 0 : 2; + 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: case CV_YUV2RGB: { - CV_Assert(scn == 3 || scn == 4); - bidx = code == CV_YUV2RGB ? 0 : 2; - dst.create(sz, CV_MAKETYPE(depth, 3)); - YUV2RGB_caller(src, dst, bidx); + if( dcn <= 0 ) + dcn = 3; + CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); + bidx = code == CV_YUV2BGR ? 0 : 2; + dst.create(sz, CV_MAKETYPE(depth, dcn)); + toRGB_caller(src, dst, bidx, "YUV2RGB"); break; } - case CV_YUV2RGB_NV12: - case CV_YUV2BGR_NV12: - case CV_YUV2RGBA_NV12: - case CV_YUV2BGRA_NV12: + case CV_YUV2RGB_NV12: case CV_YUV2BGR_NV12: + case CV_YUV2RGBA_NV12: case CV_YUV2BGRA_NV12: { CV_Assert(scn == 1); CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); - dcn = code == CV_YUV2BGRA_NV12 || code == CV_YUV2RGBA_NV12 ? 4 : 3; + dcn = code == CV_YUV2BGRA_NV12 || code == CV_YUV2RGBA_NV12 ? 4 : 3; bidx = code == CV_YUV2BGRA_NV12 || code == CV_YUV2BGR_NV12 ? 0 : 2; 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: @@ -274,20 +283,117 @@ 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: case CV_YCrCb2RGB: { + if( dcn <= 0 ) + dcn = 3; + CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); + bidx = code == CV_YCrCb2BGR ? 0 : 2; + dst.create(sz, CV_MAKETYPE(depth, dcn)); + toRGB_caller(src, dst, bidx, "YCrCb2RGB"); break; } /* case CV_BGR5652GRAY: case CV_BGR5552GRAY: case CV_GRAY2BGR565: case CV_GRAY2BGR555: - case CV_BGR2YCrCb: case CV_RGB2YCrCb: - case CV_BGR2XYZ: case CV_RGB2XYZ: - case CV_XYZ2BGR: case CV_XYZ2RGB: + */ + case CV_BGR2XYZ: + case CV_RGB2XYZ: + { + CV_Assert(scn == 3 || scn == 4); + bidx = code == CV_BGR2XYZ ? 0 : 2; + dst.create(sz, CV_MAKE_TYPE(depth, 3)); + + void * pdata = NULL; + if (depth == CV_32F) + { + float coeffs[] = + { + 0.412453f, 0.357580f, 0.180423f, + 0.212671f, 0.715160f, 0.072169f, + 0.019334f, 0.119193f, 0.950227f + }; + if (bidx == 0) + { + std::swap(coeffs[0], coeffs[2]); + std::swap(coeffs[3], coeffs[5]); + std::swap(coeffs[6], coeffs[8]); + } + pdata = coeffs; + } + else + { + int coeffs[] = + { + 1689, 1465, 739, + 871, 2929, 296, + 79, 488, 3892 + }; + if (bidx == 0) + { + std::swap(coeffs[0], coeffs[2]); + std::swap(coeffs[3], coeffs[5]); + std::swap(coeffs[6], coeffs[8]); + } + pdata = coeffs; + } + oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32FC1 : CV_32SC1, pdata); + + fromRGB_caller(src, dst, bidx, "RGB2XYZ", oclCoeffs); + break; + } + case CV_XYZ2BGR: + case CV_XYZ2RGB: + { + if (dcn <= 0) + dcn = 3; + CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); + bidx = code == CV_XYZ2BGR ? 0 : 2; + dst.create(sz, CV_MAKE_TYPE(depth, dcn)); + + void * pdata = NULL; + if (depth == CV_32F) + { + float coeffs[] = + { + 3.240479f, -1.53715f, -0.498535f, + -0.969256f, 1.875991f, 0.041556f, + 0.055648f, -0.204043f, 1.057311f + }; + if (bidx == 0) + { + std::swap(coeffs[0], coeffs[6]); + std::swap(coeffs[1], coeffs[7]); + std::swap(coeffs[2], coeffs[8]); + } + pdata = coeffs; + } + else + { + int coeffs[] = + { + 13273, -6296, -2042, + -3970, 7684, 170, + 228, -836, 4331 + }; + if (bidx == 0) + { + std::swap(coeffs[0], coeffs[6]); + std::swap(coeffs[1], coeffs[7]); + std::swap(coeffs[2], coeffs[8]); + } + pdata = coeffs; + } + oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32FC1 : CV_32SC1, pdata); + + toRGB_caller(src, dst, bidx, "XYZ2RGB", oclCoeffs); + break; + } + /* case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL: case CV_BGR2HLS: case CV_RGB2HLS: case CV_BGR2HLS_FULL: case CV_RGB2HLS_FULL: case CV_HSV2BGR: case CV_HSV2RGB: case CV_HSV2BGR_FULL: case CV_HSV2RGB_FULL: @@ -297,7 +403,6 @@ void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" ); } } -} void cv::ocl::cvtColor(const oclMat &src, oclMat &dst, int code, int dcn) { diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 01286f7ad7..210d1b766e 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -46,36 +46,31 @@ /**************************************PUBLICFUNC*************************************/ -#if defined (DOUBLE_SUPPORT) -#pragma OPENCL EXTENSION cl_khr_fp64:enable -#endif - -#if defined (DEPTH_0) +#ifdef DEPTH_0 #define DATA_TYPE uchar +#define COEFF_TYPE int #define MAX_NUM 255 #define HALF_MAX 128 -#define SAT_CAST(num) convert_uchar_sat(num) +#define SAT_CAST(num) convert_uchar_sat_rte(num) #endif -#if defined (DEPTH_2) +#ifdef DEPTH_2 #define DATA_TYPE ushort +#define COEFF_TYPE int #define MAX_NUM 65535 #define HALF_MAX 32768 -#define SAT_CAST(num) convert_ushort_sat(num) +#define SAT_CAST(num) convert_ushort_sat_rte(num) #endif -#if defined (DEPTH_5) +#ifdef DEPTH_5 #define DATA_TYPE float +#define COEFF_TYPE float #define MAX_NUM 1.0f #define HALF_MAX 0.5f #define SAT_CAST(num) (num) #endif -#ifndef DATA_TYPE - #define DATA_TYPE UNDEFINED -#endif - -#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n)) +#define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n)) enum { @@ -89,18 +84,18 @@ enum ///////////////////////////////////// RGB <-> GRAY ////////////////////////////////////// -__kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step, int channels, +__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) { - const int x = get_global_id(0); - const int y = get_global_id(1); + 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 * channels); + int src_idx = mad24(y, src_step, src_offset + (x << 2)); int dst_idx = mad24(y, dst_step, dst_offset + x); -#if defined (DEPTH_5) +#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); @@ -108,22 +103,25 @@ __kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step, int chann } } -__kernel void Gray2RGB(int cols,int rows,int src_step,int dst_step, +__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) { - const int x = get_global_id(0); - const int y = get_global_id(1); + 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 * 4); + int dst_idx = mad24(y, dst_step, dst_offset + (x << 2)); + DATA_TYPE val = src[src_idx]; - dst[dst_idx++] = val; - dst[dst_idx++] = val; - dst[dst_idx++] = val; - dst[dst_idx] = MAX_NUM; + dst[dst_idx] = val; + dst[dst_idx + 1] = val; + dst[dst_idx + 2] = val; +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; +#endif } } @@ -132,7 +130,7 @@ __kernel void Gray2RGB(int cols,int rows,int src_step,int dst_step, __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 channels, +__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) { @@ -141,35 +139,34 @@ __kernel void RGB2YUV(int cols,int rows,int src_step,int dst_step,int channels, if (y < rows && x < cols) { - x *= channels; + x <<= 2; int src_idx = mad24(y, src_step, src_offset + x); int dst_idx = mad24(y, dst_step, dst_offset + x); - dst += dst_idx; - const DATA_TYPE rgb[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]}; + DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; -#if defined (DEPTH_5) +#ifdef DEPTH_5 __constant float * coeffs = c_RGB2YUVCoeffs_f; - const DATA_TYPE Y = rgb[0] * coeffs[bidx] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx^2]; - const DATA_TYPE Cr = (rgb[bidx] - Y) * coeffs[3] + HALF_MAX; - const DATA_TYPE Cb = (rgb[bidx^2] - Y) * coeffs[4] + HALF_MAX; + 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; - const int delta = HALF_MAX * (1 << yuv_shift); - const int Y = CV_DESCALE(rgb[0] * coeffs[bidx] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx^2], yuv_shift); - const int Cr = CV_DESCALE((rgb[bidx] - Y) * coeffs[3] + delta, yuv_shift); - const int Cb = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[4] + delta, yuv_shift); + 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[0] = SAT_CAST( Y ); - dst[1] = SAT_CAST( Cr ); - dst[2] = SAT_CAST( Cb ); + 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 channels, +__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) { @@ -178,27 +175,29 @@ __kernel void YUV2RGB(int cols,int rows,int src_step,int dst_step,int channels, if (y < rows && x < cols) { - x *= channels; + x <<= 2; int src_idx = mad24(y, src_step, src_offset + x); int dst_idx = mad24(y, dst_step, dst_offset + x); - dst += dst_idx; - const DATA_TYPE yuv[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]}; + DATA_TYPE yuv[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; -#if defined (DEPTH_5) +#ifdef DEPTH_5 __constant float * coeffs = c_YUV2RGBCoeffs_f; - const float b = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[3]; - const float g = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1]; - const float r = yuv[0] + (yuv[1] - HALF_MAX) * coeffs[0]; + 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; - const int b = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[3], yuv_shift); - const int g = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1], yuv_shift); - const int r = yuv[0] + CV_DESCALE((yuv[1] - HALF_MAX) * coeffs[0], yuv_shift); + 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[bidx^2] = SAT_CAST( b ); - dst[1] = SAT_CAST( g ); - dst[bidx] = SAT_CAST( r ); + 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 } } @@ -209,17 +208,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 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); @@ -261,12 +260,12 @@ __kernel void YUV2RGBA_NV12(int cols,int rows,int src_step,int dst_step, } } -///////////////////////////////////// RGB <-> YUV ////////////////////////////////////// +///////////////////////////////////// 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 channels, +__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) { @@ -275,28 +274,273 @@ __kernel void RGB2YCrCb(int cols,int rows,int src_step,int dst_step,int channels if (y < rows && x < cols) { - x *= channels; + x <<= 2; int src_idx = mad24(y, src_step, src_offset + x); int dst_idx = mad24(y, dst_step, dst_offset + x); - dst += dst_idx; - const DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; + DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; -#if defined (DEPTH_5) +#ifdef DEPTH_5 __constant float * coeffs = c_RGB2YCrCbCoeffs_f; - const DATA_TYPE Y = rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx]; - const DATA_TYPE Cr = (rgb[bidx^2] - Y) * coeffs[3] + HALF_MAX; - const DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX; + 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; - const int delta = HALF_MAX * (1 << yuv_shift); - const int Y = CV_DESCALE(rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx], yuv_shift); - const int Cr = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[3] + delta, yuv_shift); - const int Cb = CV_DESCALE((rgb[bidx] - Y) * coeffs[4] + delta, yuv_shift); + 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[0] = SAT_CAST( Y ); - dst[1] = SAT_CAST( Cr ); - dst[2] = SAT_CAST( Cb ); + 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 }; + +__kernel void YCrCb2RGB(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 ycrcb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; + +#ifdef DEPTH_5 + __constant float * coeff = c_YCrCb2RGBCoeffs_f; + float r = ycrcb[0] + coeff[0] * (ycrcb[1] - HALF_MAX); + float g = ycrcb[0] + coeff[1] * (ycrcb[1] - HALF_MAX) + coeff[2] * (ycrcb[2] - HALF_MAX); + float b = ycrcb[0] + coeff[3] * (ycrcb[2] - HALF_MAX); +#else + __constant int * coeff = c_YCrCb2RGBCoeffs_i; + int r = ycrcb[0] + CV_DESCALE(coeff[0] * (ycrcb[1] - HALF_MAX), yuv_shift); + int g = ycrcb[0] + CV_DESCALE(coeff[1] * (ycrcb[1] - HALF_MAX) + coeff[2] * (ycrcb[2] - HALF_MAX), yuv_shift); + int b = ycrcb[0] + CV_DESCALE(coeff[3] * (ycrcb[2] - HALF_MAX), yuv_shift); +#endif + + dst[dst_idx + (bidx^2)] = SAT_CAST(r); + dst[dst_idx + 1] = SAT_CAST(g); + dst[dst_idx + bidx] = SAT_CAST(b); +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; +#endif + } +} + +///////////////////////////////////// RGB <-> XYZ ////////////////////////////////////// + +__kernel void RGB2XYZ(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, __constant COEFF_TYPE * coeffs) +{ + int dx = get_global_id(0); + int dy = get_global_id(1); + + if (dy < rows && dx < cols) + { + dx <<= 2; + int src_idx = mad24(dy, src_step, src_offset + dx); + int dst_idx = mad24(dy, dst_step, dst_offset + dx); + + DATA_TYPE r = src[src_idx], g = src[src_idx + 1], b = src[src_idx + 2]; + +#ifdef DEPTH_5 + float x = r * coeffs[0] + g * coeffs[1] + b * coeffs[2]; + float y = r * coeffs[3] + g * coeffs[4] + b * coeffs[5]; + float z = r * coeffs[6] + g * coeffs[7] + b * coeffs[8]; +#else + int x = CV_DESCALE(r * coeffs[0] + g * coeffs[1] + b * coeffs[2], xyz_shift); + int y = CV_DESCALE(r * coeffs[3] + g * coeffs[4] + b * coeffs[5], xyz_shift); + int z = CV_DESCALE(r * coeffs[6] + g * coeffs[7] + b * coeffs[8], xyz_shift); +#endif + dst[dst_idx] = SAT_CAST(x); + dst[dst_idx + 1] = SAT_CAST(y); + dst[dst_idx + 2] = SAT_CAST(z); + } +} + +__kernel void XYZ2RGB(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, __constant COEFF_TYPE * coeffs) +{ + int dx = get_global_id(0); + int dy = get_global_id(1); + + if (dy < rows && dx < cols) + { + dx <<= 2; + int src_idx = mad24(dy, src_step, src_offset + dx); + int dst_idx = mad24(dy, dst_step, dst_offset + dx); + + DATA_TYPE x = src[src_idx], y = src[src_idx + 1], z = src[src_idx + 2]; + +#ifdef DEPTH_5 + float b = x * coeffs[0] + y * coeffs[1] + z * coeffs[2]; + float g = x * coeffs[3] + y * coeffs[4] + z * coeffs[5]; + float r = x * coeffs[6] + y * coeffs[7] + z * coeffs[8]; +#else + int b = CV_DESCALE(x * coeffs[0] + y * coeffs[1] + z * coeffs[2], xyz_shift); + int g = CV_DESCALE(x * coeffs[3] + y * coeffs[4] + z * coeffs[5], xyz_shift); + int r = CV_DESCALE(x * coeffs[6] + y * coeffs[7] + z * coeffs[8], xyz_shift); +#endif + dst[dst_idx] = SAT_CAST(b); + dst[dst_idx + 1] = SAT_CAST(g); + dst[dst_idx + 2] = SAT_CAST(r); +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; +#endif + } +} + +///////////////////////////////////// RGB[A] <-> BGR[A] ////////////////////////////////////// + +__kernel void RGB(int cols, int rows, int src_step, int dst_step, + __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); + +#ifdef REVERSE + dst[dst_idx] = src[src_idx + 2]; + dst[dst_idx + 1] = src[src_idx + 1]; + dst[dst_idx + 2] = src[src_idx]; +#elif defined ORDER + dst[dst_idx] = src[src_idx]; + dst[dst_idx + 1] = src[src_idx + 1]; + dst[dst_idx + 2] = src[src_idx + 2]; +#endif + +#if dcn == 4 +#if scn == 3 + dst[dst_idx + 3] = MAX_NUM; +#else + dst[dst_idx + 3] = src[src_idx + 3]; +#endif +#endif + } +} + +///////////////////////////////////// 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 ////////////////////////////////////// + +__kernel void BGR5x52Gray(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); + int t = src[src_idx]; + +#if greenbits == 6 + dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + + ((t >> 3) & 0xfc)*G2Y + + ((t >> 8) & 0xf8)*R2Y, yuv_shift); +#else + dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + + ((t >> 2) & 0xf8)*G2Y + + ((t >> 7) & 0xf8)*R2Y, yuv_shift); +#endif + } +} + +__kernel void Gray2BGR5x5(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); + int dst_idx = mad24(y, dst_step, dst_offset + x); + int t = src[src_idx]; + +#if greenbits == 6 + dst[dst_idx] = (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8)); +#else + t >>= 3; + dst[dst_idx] = (ushort)(t|(t << 5)|(t << 10)); +#endif } } diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index 4b6a6e41b9..732a0f8e45 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -46,9 +46,8 @@ #include "test_precomp.hpp" #ifdef HAVE_OPENCL -namespace -{ using namespace testing; +using namespace cv; /////////////////////////////////////////////////////////////////////////////////////////////////////// // cvtColor @@ -59,20 +58,20 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) bool use_roi; // src mat - cv::Mat src1; - cv::Mat dst1; + Mat src; + Mat dst; // src mat with roi - cv::Mat src1_roi; - cv::Mat dst1_roi; + Mat src_roi; + Mat dst_roi; // ocl dst mat for testing - cv::ocl::oclMat gsrc1_whole; - cv::ocl::oclMat gdst1_whole; + ocl::oclMat gsrc_whole; + ocl::oclMat gdst_whole; // ocl mat with roi - cv::ocl::oclMat gsrc1_roi; - cv::ocl::oclMat gdst1_roi; + ocl::oclMat gsrc_roi; + ocl::oclMat gdst_roi; virtual void SetUp() { @@ -87,19 +86,23 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) Size roiSize = randomSize(1, MAX_VALUE); Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); - randomSubMat(src1, src1_roi, roiSize, srcBorder, srcType, 2, 100); + randomSubMat(src, src_roi, roiSize, srcBorder, srcType, 2, 100); - Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0); - randomSubMat(dst1, dst1_roi, roiSize, dst1Border, dstType, 5, 16); + Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, dstType, 5, 16); - generateOclMat(gsrc1_whole, gsrc1_roi, src1, roiSize, srcBorder); - generateOclMat(gdst1_whole, gdst1_roi, dst1, roiSize, dst1Border); + generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder); + generateOclMat(gdst_whole, gdst_roi, dst, roiSize, dstBorder); } void Near(double threshold = 1e-3) { - EXPECT_MAT_NEAR(dst1, gdst1_whole, threshold); - EXPECT_MAT_NEAR(dst1_roi, gdst1_roi, threshold); + Mat whole, roi; + gdst_whole.download(whole); + gdst_roi.download(roi); + + EXPECT_MAT_NEAR(dst, whole, threshold); + EXPECT_MAT_NEAR(dst_roi, roi, threshold); } void doTest(int channelsIn, int channelsOut, int code) @@ -108,78 +111,112 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) { random_roi(channelsIn, channelsOut); - cv::cvtColor(src1_roi, dst1_roi, code); - cv::ocl::cvtColor(gsrc1_roi, gdst1_roi, code); + cvtColor(src_roi, dst_roi, code, channelsOut); + ocl::cvtColor(gsrc_roi, gdst_roi, code, channelsOut); Near(); } } }; -#define CVTCODE(name) cv::COLOR_ ## name +#define CVTCODE(name) COLOR_ ## name -OCL_TEST_P(CvtColor, RGB2GRAY) -{ - doTest(3, 1, CVTCODE(RGB2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2RGB) -{ - doTest(1, 3, CVTCODE(GRAY2RGB)); -}; +// RGB[A] <-> BGR[A] -OCL_TEST_P(CvtColor, BGR2GRAY) -{ - doTest(3, 1, CVTCODE(BGR2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2BGR) -{ - doTest(1, 3, CVTCODE(GRAY2BGR)); -}; +OCL_TEST_P(CvtColor, BGR2BGRA) { doTest(3, 4, CVTCODE(BGR2BGRA)); } +OCL_TEST_P(CvtColor, RGB2RGBA) { doTest(3, 4, CVTCODE(BGR2BGRA)); } +OCL_TEST_P(CvtColor, BGRA2BGR) { doTest(4, 3, CVTCODE(BGRA2BGR)); } +OCL_TEST_P(CvtColor, RGBA2RGB) { doTest(4, 3, CVTCODE(BGRA2BGR)); } +OCL_TEST_P(CvtColor, BGR2RGBA) { doTest(3, 4, CVTCODE(BGR2RGBA)); } +OCL_TEST_P(CvtColor, RGB2BGRA) { doTest(3, 4, CVTCODE(BGR2RGBA)); } +OCL_TEST_P(CvtColor, RGBA2BGR) { doTest(4, 3, CVTCODE(RGBA2BGR)); } +OCL_TEST_P(CvtColor, BGRA2RGB) { doTest(4, 3, CVTCODE(RGBA2BGR)); } +OCL_TEST_P(CvtColor, BGR2RGB) { doTest(3, 3, CVTCODE(BGR2RGB)); } +OCL_TEST_P(CvtColor, RGB2BGR) { doTest(3, 3, CVTCODE(BGR2RGB)); } +OCL_TEST_P(CvtColor, BGRA2RGBA) { doTest(4, 4, CVTCODE(BGRA2RGBA)); } +OCL_TEST_P(CvtColor, RGBA2BGRA) { doTest(4, 4, CVTCODE(BGRA2RGBA)); } -OCL_TEST_P(CvtColor, RGBA2GRAY) -{ - doTest(3, 1, CVTCODE(RGBA2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2RGBA) -{ - doTest(1, 3, CVTCODE(GRAY2RGBA)); -}; +// RGB <-> Gray -OCL_TEST_P(CvtColor, BGRA2GRAY) -{ - doTest(3, 1, CVTCODE(BGRA2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2BGRA) -{ - doTest(1, 3, CVTCODE(GRAY2BGRA)); -}; +OCL_TEST_P(CvtColor, RGB2GRAY) { doTest(3, 1, CVTCODE(RGB2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2RGB) { doTest(1, 3, CVTCODE(GRAY2RGB)); } +OCL_TEST_P(CvtColor, BGR2GRAY) { doTest(3, 1, CVTCODE(BGR2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2BGR) { doTest(1, 3, CVTCODE(GRAY2BGR)); } +OCL_TEST_P(CvtColor, RGBA2GRAY) { doTest(4, 1, CVTCODE(RGBA2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2RGBA) { doTest(1, 4, CVTCODE(GRAY2RGBA)); } +OCL_TEST_P(CvtColor, BGRA2GRAY) { doTest(4, 1, CVTCODE(BGRA2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2BGRA) { doTest(1, 4, CVTCODE(GRAY2BGRA)); } -OCL_TEST_P(CvtColor, RGB2YUV) -{ - doTest(3, 3, CVTCODE(RGB2YUV)); -} -OCL_TEST_P(CvtColor, BGR2YUV) -{ - doTest(3, 3, CVTCODE(BGR2YUV)); -} -OCL_TEST_P(CvtColor, YUV2RGB) -{ - doTest(3, 3, CVTCODE(YUV2RGB)); -} -OCL_TEST_P(CvtColor, YUV2BGR) -{ - doTest(3, 3, CVTCODE(YUV2BGR)); -} -OCL_TEST_P(CvtColor, RGB2YCrCb) -{ - doTest(3, 3, CVTCODE(RGB2YCrCb)); -} -OCL_TEST_P(CvtColor, BGR2YCrCb) -{ - doTest(3, 3, CVTCODE(BGR2YCrCb)); -} +// RGB <-> YUV -struct CvtColor_YUV420 : CvtColor +OCL_TEST_P(CvtColor, RGB2YUV) { doTest(3, 3, CVTCODE(RGB2YUV)); } +OCL_TEST_P(CvtColor, BGR2YUV) { doTest(3, 3, CVTCODE(BGR2YUV)); } +OCL_TEST_P(CvtColor, RGBA2YUV) { doTest(4, 3, CVTCODE(RGB2YUV)); } +OCL_TEST_P(CvtColor, BGRA2YUV) { doTest(4, 3, CVTCODE(BGR2YUV)); } +OCL_TEST_P(CvtColor, YUV2RGB) { doTest(3, 3, CVTCODE(YUV2RGB)); } +OCL_TEST_P(CvtColor, YUV2BGR) { doTest(3, 3, CVTCODE(YUV2BGR)); } +OCL_TEST_P(CvtColor, YUV2RGBA) { doTest(3, 4, CVTCODE(YUV2RGB)); } +OCL_TEST_P(CvtColor, YUV2BGRA) { doTest(3, 4, CVTCODE(YUV2BGR)); } + +// RGB <-> YCrCb + +OCL_TEST_P(CvtColor, RGB2YCrCb) { doTest(3, 3, CVTCODE(RGB2YCrCb)); } +OCL_TEST_P(CvtColor, BGR2YCrCb) { doTest(3, 3, CVTCODE(BGR2YCrCb)); } +OCL_TEST_P(CvtColor, RGBA2YCrCb) { doTest(4, 3, CVTCODE(RGB2YCrCb)); } +OCL_TEST_P(CvtColor, BGRA2YCrCb) { doTest(4, 3, CVTCODE(BGR2YCrCb)); } +OCL_TEST_P(CvtColor, YCrCb2RGB) { doTest(3, 3, CVTCODE(YCrCb2RGB)); } +OCL_TEST_P(CvtColor, YCrCb2BGR) { doTest(3, 3, CVTCODE(YCrCb2BGR)); } +OCL_TEST_P(CvtColor, YCrCb2RGBA) { doTest(3, 4, CVTCODE(YCrCb2RGB)); } +OCL_TEST_P(CvtColor, YCrCb2BGRA) { doTest(3, 4, CVTCODE(YCrCb2BGR)); } + +// RGB <-> XYZ + +OCL_TEST_P(CvtColor, RGB2XYZ) { doTest(3, 3, CVTCODE(RGB2XYZ)); } +OCL_TEST_P(CvtColor, BGR2XYZ) { doTest(3, 3, CVTCODE(BGR2XYZ)); } +OCL_TEST_P(CvtColor, RGBA2XYZ) { doTest(4, 3, CVTCODE(RGB2XYZ)); } +OCL_TEST_P(CvtColor, BGRA2XYZ) { doTest(4, 3, CVTCODE(BGR2XYZ)); } + +OCL_TEST_P(CvtColor, XYZ2RGB) { doTest(3, 3, CVTCODE(XYZ2RGB)); } +OCL_TEST_P(CvtColor, XYZ2BGR) { doTest(3, 3, CVTCODE(XYZ2BGR)); } +OCL_TEST_P(CvtColor, XYZ2RGBA) { doTest(3, 4, CVTCODE(XYZ2RGB)); } +OCL_TEST_P(CvtColor, XYZ2BGRA) { doTest(3, 4, CVTCODE(XYZ2BGR)); } + +// RGB5x5 <-> RGB + +typedef CvtColor CvtColor8u; + +OCL_TEST_P(CvtColor8u, BGR5652BGR) { doTest(2, 3, CVTCODE(BGR5652BGR)); } +OCL_TEST_P(CvtColor8u, BGR5652RGB) { doTest(2, 3, CVTCODE(BGR5652RGB)); } +OCL_TEST_P(CvtColor8u, BGR5652BGRA) { doTest(2, 4, CVTCODE(BGR5652BGRA)); } +OCL_TEST_P(CvtColor8u, BGR5652RGBA) { doTest(2, 4, CVTCODE(BGR5652RGBA)); } + +OCL_TEST_P(CvtColor8u, BGR5552BGR) { doTest(2, 3, CVTCODE(BGR5552BGR)); } +OCL_TEST_P(CvtColor8u, BGR5552RGB) { doTest(2, 3, CVTCODE(BGR5552RGB)); } +OCL_TEST_P(CvtColor8u, BGR5552BGRA) { doTest(2, 4, CVTCODE(BGR5552BGRA)); } +OCL_TEST_P(CvtColor8u, BGR5552RGBA) { doTest(2, 4, CVTCODE(BGR5552RGBA)); } + +OCL_TEST_P(CvtColor8u, BGR2BGR565) { doTest(3, 2, CVTCODE(BGR2BGR565)); } +OCL_TEST_P(CvtColor8u, RGB2BGR565) { doTest(3, 2, CVTCODE(RGB2BGR565)); } +OCL_TEST_P(CvtColor8u, BGRA2BGR565) { doTest(4, 2, CVTCODE(BGRA2BGR565)); } +OCL_TEST_P(CvtColor8u, RGBA2BGR565) { doTest(4, 2, CVTCODE(RGBA2BGR565)); } + +OCL_TEST_P(CvtColor8u, BGR2BGR555) { doTest(3, 2, CVTCODE(BGR2BGR555)); } +OCL_TEST_P(CvtColor8u, RGB2BGR555) { doTest(3, 2, CVTCODE(RGB2BGR555)); } +OCL_TEST_P(CvtColor8u, BGRA2BGR555) { doTest(4, 2, CVTCODE(BGRA2BGR555)); } +OCL_TEST_P(CvtColor8u, RGBA2BGR555) { doTest(4, 2, CVTCODE(RGBA2BGR555)); } + +// RGB5x5 <-> Gray + +OCL_TEST_P(CvtColor8u, BGR5652GRAY) { doTest(2, 1, CVTCODE(BGR5652GRAY)); } +OCL_TEST_P(CvtColor8u, BGR5552GRAY) { doTest(2, 1, CVTCODE(BGR5552GRAY)); } + +OCL_TEST_P(CvtColor8u, GRAY2BGR565) { doTest(1, 2, CVTCODE(GRAY2BGR565)); } +OCL_TEST_P(CvtColor8u, GRAY2BGR555) { doTest(1, 2, CVTCODE(GRAY2BGR555)); } + +// YUV -> RGBA_NV12 + +struct CvtColor_YUV420 : + public CvtColor { void random_roi(int channelsIn, int channelsOut) { @@ -190,50 +227,33 @@ struct CvtColor_YUV420 : CvtColor roiSize.width *= 2; roiSize.height *= 3; Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); - randomSubMat(src1, src1_roi, roiSize, srcBorder, srcType, 2, 100); + randomSubMat(src, src_roi, roiSize, srcBorder, srcType, 2, 100); - Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0); - randomSubMat(dst1, dst1_roi, roiSize, dst1Border, dstType, 5, 16); + Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, dstType, 5, 16); - generateOclMat(gsrc1_whole, gsrc1_roi, src1, roiSize, srcBorder); - generateOclMat(gdst1_whole, gdst1_roi, dst1, roiSize, dst1Border); + generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder); + generateOclMat(gdst_whole, gdst_roi, dst, roiSize, dstBorder); } }; -OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) -{ - doTest(1, 4, CV_YUV2RGBA_NV12); -}; +OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) { doTest(1, 4, CV_YUV2RGBA_NV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) { doTest(1, 4, CV_YUV2BGRA_NV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { doTest(1, 3, CV_YUV2RGB_NV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { doTest(1, 3, CV_YUV2BGR_NV12); } -OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) -{ - doTest(1, 4, CV_YUV2BGRA_NV12); -}; - -OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) -{ - doTest(1, 3, CV_YUV2RGB_NV12); -}; - -OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) -{ - doTest(1, 3, CV_YUV2BGR_NV12); -}; +INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor8u, + testing::Combine(testing::Values(MatDepth(CV_8U)), Bool())); INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor, testing::Combine( testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)), - Bool() - ) - ); + Bool())); INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor_YUV420, testing::Combine( testing::Values(MatDepth(CV_8U)), - Bool() - ) - ); + Bool())); -} #endif