From 7452eef6e972cc787d858bca52913fd9d377c693 Mon Sep 17 00:00:00 2001 From: Chuanbo Weng Date: Mon, 20 Oct 2014 11:50:46 +0800 Subject: [PATCH 01/26] Correctly enable OpenCL mode in tapi's hog example. For current OpenCV-CL architecture, if the data buffer allocated in UMat are cpu buffer(not ocl buffer) under cpu mode, and then pass this UMat to an OpenCL kernel as an argument, the OpenCL path will fail and fallback to cpu mode. Take HOGDescriptor::oclSvmDetector as an example: ocl::setUseOpenCL(false); //data allocated in hog.oclSvmDetector will be cpu buffer hog.setSVMDetector(HOGDescriptor::getDaimlerPeopleDetector()); ocl::setUseOpenCL(true); //We enabled OpenCL, but hog.oclSvmDetector are cpu buffer, //so it will fail in the function ocl_classify_hists //when reach to this line //idx = k.set(idx, ocl::KernelArg::PtrReadOnly(detector)); hog.detectMultiScale(img, found, hit_threshold, win_stride, Size(0, 0), scale, gr_threshold); Similar problems heppen on img_aux and img. So we should re-define or re-set these UMat when do mode switch (CPU -> OpenCL) in order to make their data be allocated by ocl and then OpenCL path will succeed. Signed-off-by: Chuanbo Weng --- samples/tapi/hog.cpp | 26 ++++++++++++++++++++------ 1 file changed, 20 insertions(+), 6 deletions(-) diff --git a/samples/tapi/hog.cpp b/samples/tapi/hog.cpp index 389e1e5bef..62d7aa44b7 100644 --- a/samples/tapi/hog.cpp +++ b/samples/tapi/hog.cpp @@ -44,6 +44,8 @@ private: //Args args; bool running; bool make_gray; + bool use_ocl; + bool ocl_switch; double scale; double resize_scale; int win_width; @@ -134,6 +136,9 @@ App::App(CommandLineParser& cmd) gamma_corr = true; write_once = false; + use_ocl = ocl::useOpenCL(); + ocl_switch = true; + cout << "Group threshold: " << gr_threshold << endl; cout << "Levels number: " << nlevels << endl; cout << "Win width: " << win_width << endl; @@ -155,7 +160,6 @@ void App::run() HOGDescriptor hog(win_size, Size(16, 16), Size(8, 8), Size(8, 8), 9, 1, -1, HOGDescriptor::L2Hys, 0.2, gamma_corr, cv::HOGDescriptor::DEFAULT_NLEVELS); - hog.setSVMDetector( HOGDescriptor::getDaimlerPeopleDetector() ); while (running) { @@ -187,13 +191,17 @@ void App::run() throw runtime_error(string("can't open image file: " + img_source)); } - UMat img_aux, img; Mat img_to_show; // Iterate over all frames while (running && !frame.empty()) { workBegin(); + if(ocl_switch){ + hog.setSVMDetector( HOGDescriptor::getDaimlerPeopleDetector() ); + ocl_switch = false; + } + UMat img_aux, img; // Change format of the image if (make_gray) cvtColor(frame, img_aux, COLOR_BGR2GRAY ); @@ -213,8 +221,12 @@ void App::run() // Perform HOG classification hogWorkBegin(); - hog.detectMultiScale(img.getMat(ACCESS_READ), found, hit_threshold, win_stride, - Size(0, 0), scale, gr_threshold); + if(use_ocl) + hog.detectMultiScale(img, found, hit_threshold, win_stride, + Size(0, 0), scale, gr_threshold); + else + hog.detectMultiScale(img.getMat(ACCESS_READ), found, hit_threshold, win_stride, + Size(0, 0), scale, gr_threshold); hogWorkEnd(); @@ -225,7 +237,7 @@ void App::run() rectangle(img_to_show, r.tl(), r.br(), Scalar(0, 255, 0), 3); } - putText(img_to_show, "Mode: CPU", Point(5, 25), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); + putText(img_to_show, use_ocl ? "Mode: OpenCL" : "Mode: CPU", Point(5, 25), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); putText(img_to_show, "FPS (HOG only): " + hogWorkFps(), Point(5, 65), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); putText(img_to_show, "FPS (total): " + workFps(), Point(5, 105), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); imshow("opencv_hog", img_to_show); @@ -272,7 +284,9 @@ void App::handleKey(char key) case 'm': case 'M': ocl::setUseOpenCL(!cv::ocl::useOpenCL()); - cout << "Switched to " << (ocl::useOpenCL() ? "OpenCL enabled" : "CPU") << " mode\n"; + ocl_switch = true; + use_ocl = ocl::useOpenCL(); + cout << "Switched to " << (use_ocl ? "OpenCL enabled" : "CPU") << " mode\n"; break; case 'g': case 'G': From 2b52bb092aabdbeb8c49ac0db1227c93558b30a9 Mon Sep 17 00:00:00 2001 From: Chuanbo Weng Date: Mon, 20 Oct 2014 18:14:38 +0800 Subject: [PATCH 02/26] Update hog.cpp Update according to vbystricky's comments --- samples/tapi/hog.cpp | 26 ++++++-------------------- 1 file changed, 6 insertions(+), 20 deletions(-) diff --git a/samples/tapi/hog.cpp b/samples/tapi/hog.cpp index 62d7aa44b7..3ccc41b185 100644 --- a/samples/tapi/hog.cpp +++ b/samples/tapi/hog.cpp @@ -44,8 +44,6 @@ private: //Args args; bool running; bool make_gray; - bool use_ocl; - bool ocl_switch; double scale; double resize_scale; int win_width; @@ -136,9 +134,6 @@ App::App(CommandLineParser& cmd) gamma_corr = true; write_once = false; - use_ocl = ocl::useOpenCL(); - ocl_switch = true; - cout << "Group threshold: " << gr_threshold << endl; cout << "Levels number: " << nlevels << endl; cout << "Win width: " << win_width << endl; @@ -160,6 +155,7 @@ void App::run() HOGDescriptor hog(win_size, Size(16, 16), Size(8, 8), Size(8, 8), 9, 1, -1, HOGDescriptor::L2Hys, 0.2, gamma_corr, cv::HOGDescriptor::DEFAULT_NLEVELS); + hog.setSVMDetector( HOGDescriptor::getDaimlerPeopleDetector() ); while (running) { @@ -191,17 +187,13 @@ void App::run() throw runtime_error(string("can't open image file: " + img_source)); } + UMat img_aux, img; Mat img_to_show; // Iterate over all frames while (running && !frame.empty()) { workBegin(); - if(ocl_switch){ - hog.setSVMDetector( HOGDescriptor::getDaimlerPeopleDetector() ); - ocl_switch = false; - } - UMat img_aux, img; // Change format of the image if (make_gray) cvtColor(frame, img_aux, COLOR_BGR2GRAY ); @@ -221,12 +213,8 @@ void App::run() // Perform HOG classification hogWorkBegin(); - if(use_ocl) - hog.detectMultiScale(img, found, hit_threshold, win_stride, - Size(0, 0), scale, gr_threshold); - else - hog.detectMultiScale(img.getMat(ACCESS_READ), found, hit_threshold, win_stride, - Size(0, 0), scale, gr_threshold); + hog.detectMultiScale(img, found, hit_threshold, win_stride, + Size(0, 0), scale, gr_threshold); hogWorkEnd(); @@ -237,7 +225,7 @@ void App::run() rectangle(img_to_show, r.tl(), r.br(), Scalar(0, 255, 0), 3); } - putText(img_to_show, use_ocl ? "Mode: OpenCL" : "Mode: CPU", Point(5, 25), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); + putText(img_to_show, ocl::useOpenCL() ? "Mode: OpenCL" : "Mode: CPU", Point(5, 25), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); putText(img_to_show, "FPS (HOG only): " + hogWorkFps(), Point(5, 65), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); putText(img_to_show, "FPS (total): " + workFps(), Point(5, 105), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); imshow("opencv_hog", img_to_show); @@ -284,9 +272,7 @@ void App::handleKey(char key) case 'm': case 'M': ocl::setUseOpenCL(!cv::ocl::useOpenCL()); - ocl_switch = true; - use_ocl = ocl::useOpenCL(); - cout << "Switched to " << (use_ocl ? "OpenCL enabled" : "CPU") << " mode\n"; + cout << "Switched to " << (ocl::useOpenCL() ? "OpenCL enabled" : "CPU") << " mode\n"; break; case 'g': case 'G': From 8ab6852927a94b4b81202e8ce4559ba8e253d653 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Tue, 21 Oct 2014 15:23:19 +0400 Subject: [PATCH 03/26] added performance validation check --- modules/ts/include/opencv2/ts/ts_perf.hpp | 2 + modules/ts/src/ts_perf.cpp | 206 +++++++++++++++++++++- 2 files changed, 203 insertions(+), 5 deletions(-) diff --git a/modules/ts/include/opencv2/ts/ts_perf.hpp b/modules/ts/include/opencv2/ts/ts_perf.hpp index 76fecba5b9..aac8ae5025 100644 --- a/modules/ts/include/opencv2/ts/ts_perf.hpp +++ b/modules/ts/include/opencv2/ts/ts_perf.hpp @@ -419,9 +419,11 @@ private: static int64 timeLimitDefault; static unsigned int iterationsLimitDefault; + unsigned int minIters; unsigned int nIters; unsigned int currentIter; unsigned int runsPerIteration; + unsigned int perfValidationStage; performance_metrics metrics; void validateMetrics(); diff --git a/modules/ts/src/ts_perf.cpp b/modules/ts/src/ts_perf.cpp index c970101e53..f2eae265ac 100644 --- a/modules/ts/src/ts_perf.cpp +++ b/modules/ts/src/ts_perf.cpp @@ -1,5 +1,16 @@ #include "precomp.hpp" +#include +#include +#include + +#if defined WIN32 || defined _WIN32 || defined WIN64 || defined _WIN64 +#ifndef NOMINMAX +#define NOMINMAX +#endif +#include +#endif + #ifdef HAVE_CUDA #include "opencv2/core/cuda.hpp" #endif @@ -35,11 +46,11 @@ static bool param_verify_sanity; static bool param_collect_impl; #endif extern bool test_ipp_check; + #ifdef HAVE_CUDA static int param_cuda_device; #endif - #ifdef ANDROID static int param_affinity_mask; static bool log_power_checkpoints; @@ -59,6 +70,8 @@ static void setCurrentThreadAffinityMask(int mask) } #endif +static double perf_stability_criteria = 0.03; // 3% + namespace { class PerfEnvironment: public ::testing::Environment @@ -635,6 +648,82 @@ void performance_metrics::clear() terminationReason = TERM_UNKNOWN; } +/*****************************************************************************************\ +* Performance validation results +\*****************************************************************************************/ + +static bool perf_validation_enabled = false; + +static std::string perf_validation_results_directory; +static std::map perf_validation_results; +static std::string perf_validation_results_outfile; + +static double perf_validation_criteria = 0.03; // 3 % +static double perf_validation_time_threshold_ms = 0.1; +static int perf_validation_idle_delay_ms = 3000; // 3 sec + +static void loadPerfValidationResults(const std::string& fileName) +{ + perf_validation_results.clear(); + std::ifstream infile(fileName.c_str()); + while (!infile.eof()) + { + std::string name; + float value = 0; + if (!(infile >> value)) + { + if (infile.eof()) + break; // it is OK + std::cout << "ERROR: Can't load performance validation results from " << fileName << "!" << std::endl; + return; + } + infile.ignore(1); + if (!(std::getline(infile, name))) + { + std::cout << "ERROR: Can't load performance validation results from " << fileName << "!" << std::endl; + return; + } + if (!name.empty() && name[name.size() - 1] == '\r') // CRLF processing on Linux + name.resize(name.size() - 1); + perf_validation_results[name] = value; + } + std::cout << "Performance validation results loaded from " << fileName << " (" << perf_validation_results.size() << " entries)" << std::endl; +} + +static void savePerfValidationResult(const std::string& name, float value) +{ + perf_validation_results[name] = value; +} + +static void savePerfValidationResults() +{ + if (!perf_validation_results_outfile.empty()) + { + std::ofstream outfile((perf_validation_results_directory + perf_validation_results_outfile).c_str()); + std::map::const_iterator i; + for (i = perf_validation_results.begin(); i != perf_validation_results.end(); ++i) + { + outfile << i->second << ';'; + outfile << i->first << std::endl; + } + outfile.close(); + std::cout << "Performance validation results saved (" << perf_validation_results.size() << " entries)" << std::endl; + } +} + +class PerfValidationEnvironment : public ::testing::Environment +{ +public: + virtual ~PerfValidationEnvironment() {} + virtual void SetUp() {} + + virtual void TearDown() + { + savePerfValidationResults(); + } +}; + + /*****************************************************************************************\ * ::perf::TestBase @@ -666,6 +755,8 @@ void TestBase::Init(const std::vector & availableImpls, "{ perf_list_impls |false |list available implementation variants and exit}" "{ perf_run_cpu |false |deprecated, equivalent to --perf_impl=plain}" "{ perf_strategy |default |specifies performance measuring strategy: default, base or simple (weak restrictions)}" + "{ perf_read_validation_results | |specifies file name with performance results from previous run}" + "{ perf_write_validation_results | |specifies file name to write performance validation results}" #ifdef ANDROID "{ perf_time_limit |6.0 |default time limit for a single test (in seconds)}" "{ perf_affinity_mask |0 |set affinity mask for the main thread}" @@ -789,6 +880,26 @@ void TestBase::Init(const std::vector & availableImpls, } #endif + { + const char* path = getenv("OPENCV_PERF_VALIDATION_DIR"); + if (path) + perf_validation_results_directory = path; + } + + std::string fileName_perf_validation_results_src = args.get("perf_read_validation_results"); + if (!fileName_perf_validation_results_src.empty()) + { + perf_validation_enabled = true; + loadPerfValidationResults(perf_validation_results_directory + fileName_perf_validation_results_src); + } + + perf_validation_results_outfile = args.get("perf_write_validation_results"); + if (!perf_validation_results_outfile.empty()) + { + perf_validation_enabled = true; + ::testing::AddGlobalTestEnvironment(new PerfValidationEnvironment()); + } + if (!args.check()) { args.printErrors(); @@ -878,7 +989,9 @@ TestBase::TestBase(): testStrategy(PERF_STRATEGY_DEFAULT), declare(this) { lastTime = totalTime = timeLimit = 0; nIters = currentIter = runsPerIteration = 0; + minIters = param_min_samples; verified = false; + perfValidationStage = 0; } #ifdef _MSC_VER # pragma warning(pop) @@ -1004,7 +1117,7 @@ bool TestBase::next() has_next = false; break; } - if (currentIter < param_min_samples) + if (currentIter < minIters) { has_next = true; break; @@ -1012,14 +1125,96 @@ bool TestBase::next() calcMetrics(); - double criteria = 0.03; // 3% if (fabs(metrics.mean) > 1e-6) - has_next = metrics.stddev > criteria * fabs(metrics.mean); + has_next = metrics.stddev > perf_stability_criteria * fabs(metrics.mean); else has_next = true; } } while (false); + if (perf_validation_enabled && !has_next) + { + calcMetrics(); + double median_ms = metrics.median * 1000.0f / metrics.frequency; + + const ::testing::TestInfo* const test_info = ::testing::UnitTest::GetInstance()->current_test_info(); + std::string name = (test_info == 0) ? "" : + std::string(test_info->test_case_name()) + "--" + test_info->name(); + + if (!perf_validation_results.empty() && !name.empty()) + { + std::map::iterator i = perf_validation_results.find(name); + bool isSame = false; + bool found = false; + bool grow = false; + if (i != perf_validation_results.end()) + { + found = true; + double prev_result = i->second; + grow = median_ms > prev_result; + isSame = fabs(median_ms - prev_result) <= perf_validation_criteria * fabs(median_ms); + if (!isSame) + { + if (perfValidationStage == 0) + { + printf("Performance is changed (samples = %d, median):\n %.2f ms (current)\n %.2f ms (previous)\n", (int)times.size(), median_ms, prev_result); + } + } + } + else + { + if (perfValidationStage == 0) + printf("New performance result is detected\n"); + } + if (!isSame) + { + if (perfValidationStage < 2) + { + if (perfValidationStage == 0 && currentIter <= minIters * 3 && currentIter < nIters) + { + unsigned int new_minIters = std::max(minIters * 5, currentIter * 3); + printf("Increase minIters from %u to %u\n", minIters, new_minIters); + minIters = new_minIters; + has_next = true; + perfValidationStage++; + } + else if (found && currentIter >= nIters && + median_ms > perf_validation_time_threshold_ms && + (grow || metrics.stddev > perf_stability_criteria * fabs(metrics.mean))) + { + printf("Performance is unstable, it may be a result of overheat problems\n"); + printf("Idle delay for %d ms... \n", perf_validation_idle_delay_ms); +#if defined WIN32 || defined _WIN32 || defined WIN64 || defined _WIN64 + Sleep(perf_validation_idle_delay_ms); +#else + usleep(perf_validation_idle_delay_ms * 1000); +#endif + has_next = true; + minIters = std::min(minIters * 5, nIters); + // reset collected samples + currentIter = 0; + times.clear(); + metrics.clear(); + perfValidationStage += 2; + } + if (!has_next) + { + printf("Assume that current result is valid\n"); + } + } + else + { + printf("Re-measured performance result: %.2f ms\n", median_ms); + } + } + } + + if (!has_next && !name.empty()) + { + savePerfValidationResult(name, (float)median_ms); + } + } + #ifdef ANDROID if (log_power_checkpoints) { @@ -1223,9 +1418,10 @@ void TestBase::validateMetrics() else if (getCurrentPerformanceStrategy() == PERF_STRATEGY_SIMPLE) { double mean = metrics.mean * 1000.0f / metrics.frequency; + double median = metrics.median * 1000.0f / metrics.frequency; double stddev = metrics.stddev * 1000.0f / metrics.frequency; double percents = stddev / mean * 100.f; - printf("[ PERFSTAT ] (samples = %d, mean = %.2f, stddev = %.2f (%.1f%%))\n", (int)metrics.samples, mean, stddev, percents); + printf("[ PERFSTAT ] (samples = %d, mean = %.2f, median = %.2f, stddev = %.2f (%.1f%%))\n", (int)metrics.samples, mean, median, stddev, percents); } else { From 85b60ee3cb109fc35a5f5f9b1c7bd74f792ff93a Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Wed, 1 Oct 2014 18:08:07 +0400 Subject: [PATCH 04/26] Added support for YUV2RGB[A]_NV21 and YUV2BGR[A]_NV21 conversion --- modules/imgproc/src/color.cpp | 18 +++++++++++------- modules/imgproc/src/opencl/cvtcolor.cl | 10 +++++++--- modules/imgproc/test/ocl/test_color.cpp | 6 +++++- 3 files changed, 23 insertions(+), 11 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 21b0b25f33..c7743cb93f 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -4848,7 +4848,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) bool ok = false; UMat src = _src.getUMat(), dst; Size sz = src.size(), dstSz = sz; - int scn = src.channels(), depth = src.depth(), bidx; + int scn = src.channels(), depth = src.depth(), bidx, uidx; int dims = 2, stripeSize = 1; ocl::Kernel k; @@ -4960,17 +4960,21 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) opts + format("-D dcn=%d -D bidx=%d", 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_YUV2RGB_NV21: case COLOR_YUV2BGR_NV21: + case COLOR_YUV2RGBA_NV12: case COLOR_YUV2BGRA_NV12: case COLOR_YUV2RGBA_NV21: case COLOR_YUV2BGRA_NV21: { CV_Assert( scn == 1 ); CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); - dcn = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2RGBA_NV12 ? 4 : 3; - bidx = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2BGR_NV12 ? 0 : 2; + dcn = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2RGBA_NV12 || + code == COLOR_YUV2BGRA_NV21 || code == COLOR_YUV2RGBA_NV21 ? 4 : 3; + bidx = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2BGR_NV12 || + code == COLOR_YUV2BGRA_NV21 || code == COLOR_YUV2BGR_NV21 ? 0 : 2; + uidx = code == COLOR_YUV2RGBA_NV21 || code == COLOR_YUV2RGB_NV21 || + code == COLOR_YUV2BGRA_NV21 || code == COLOR_YUV2BGR_NV21 ? 1 : 0; dstSz = Size(sz.width, sz.height * 2 / 3); - k.create("YUV2RGB_NV12", ocl::imgproc::cvtcolor_oclsrc, - opts + format("-D dcn=%d -D bidx=%d", dcn, bidx)); + k.create("YUV2RGB_NVx", ocl::imgproc::cvtcolor_oclsrc, + opts + format("-D dcn=%d -D bidx=%d -D uidx=%d", dcn, bidx, uidx)); break; } case COLOR_BGR2YCrCb: diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 365f470c11..b647f80045 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -111,6 +111,10 @@ enum #define B_COMP w #endif +#ifndef uidx +#define uidx 0 +#endif + #define __CAT(x, y) x##y #define CAT(x, y) __CAT(x, y) @@ -297,7 +301,7 @@ __constant int ITUR_BT_601_CVG = 852492; __constant int ITUR_BT_601_CVR = 1673527; __constant int ITUR_BT_601_SHIFT = 20; -__kernel void YUV2RGB_NV12(__global const uchar* srcptr, int src_step, int src_offset, +__kernel void YUV2RGB_NVx(__global const uchar* srcptr, int src_step, int src_offset, __global uchar* dstptr, int dst_step, int dt_offset, int rows, int cols) { @@ -321,8 +325,8 @@ __kernel void YUV2RGB_NV12(__global const uchar* srcptr, int src_step, int src_o int Y3 = ysrc[src_step]; int Y4 = ysrc[src_step + 1]; - int U = usrc[0] - 128; - int V = usrc[1] - 128; + int U = usrc[uidx] - 128; + int V = usrc[1 - uidx] - 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; diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index 818d6a85ab..2130eca20e 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -320,7 +320,7 @@ OCL_TEST_P(CvtColor8u32f, Luv2RGBA) { performTest(3, 4, CVTCODE(Luv2RGB), depth OCL_TEST_P(CvtColor8u32f, Luv2LBGRA) { performTest(3, 4, CVTCODE(Luv2LBGR), depth == CV_8U ? 1 : 1e-5); } OCL_TEST_P(CvtColor8u32f, Luv2LRGBA) { performTest(3, 4, CVTCODE(Luv2LRGB), depth == CV_8U ? 1 : 1e-5); } -// YUV -> RGBA_NV12 +// YUV -> RGBA_NVx struct CvtColor_YUV420 : public CvtColor @@ -348,6 +348,10 @@ OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) { performTest(1, 4, COLOR_YUV2RGBA_NV OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) { performTest(1, 4, COLOR_YUV2BGRA_NV12); } OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { performTest(1, 3, COLOR_YUV2RGB_NV12); } OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { performTest(1, 3, COLOR_YUV2BGR_NV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV21) { performTest(1, 4, COLOR_YUV2RGBA_NV21); } +OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV21) { performTest(1, 4, COLOR_YUV2BGRA_NV21); } +OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV21) { performTest(1, 3, COLOR_YUV2RGB_NV21); } +OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV21) { performTest(1, 3, COLOR_YUV2BGR_NV21); } OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u, From 1cc17a7186017b0e6178f0a78ba6c7da820edfe8 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Thu, 2 Oct 2014 18:10:14 +0400 Subject: [PATCH 05/26] Added OCL code for YUV2BGR_YV12 and YUV2BGR_IYUV color conversions --- modules/imgproc/src/color.cpp | 22 ++++++- modules/imgproc/src/opencl/cvtcolor.cl | 85 +++++++++++++++++++++++-- modules/imgproc/test/ocl/test_color.cpp | 10 ++- 3 files changed, 111 insertions(+), 6 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index c7743cb93f..324b824ae0 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -4973,10 +4973,30 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) code == COLOR_YUV2BGRA_NV21 || code == COLOR_YUV2BGR_NV21 ? 1 : 0; dstSz = Size(sz.width, sz.height * 2 / 3); - k.create("YUV2RGB_NVx", ocl::imgproc::cvtcolor_oclsrc, + globalsize[0] = dstSz.width / 2; globalsize[1] = (dstSz.height/2 + pxPerWIy - 1) / pxPerWIy; + k.create("YUV2RGB_NV", ocl::imgproc::cvtcolor_oclsrc, opts + format("-D dcn=%d -D bidx=%d -D uidx=%d", dcn, bidx, uidx)); break; } + case COLOR_YUV2BGR_YV12: case COLOR_YUV2RGB_YV12: case COLOR_YUV2BGRA_YV12: case COLOR_YUV2RGBA_YV12: + case COLOR_YUV2BGR_IYUV: case COLOR_YUV2RGB_IYUV: case COLOR_YUV2BGRA_IYUV: case COLOR_YUV2RGBA_IYUV: + { + CV_Assert( scn == 1 ); + CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); + dcn = code == COLOR_YUV2BGRA_YV12 || code == COLOR_YUV2RGBA_YV12 || + code == COLOR_YUV2BGRA_IYUV || code == COLOR_YUV2RGBA_IYUV ? 4 : 3; + bidx = code == COLOR_YUV2BGRA_YV12 || code == COLOR_YUV2BGR_YV12 || + code == COLOR_YUV2BGRA_IYUV || code == COLOR_YUV2BGR_IYUV ? 0 : 2; + uidx = code == COLOR_YUV2BGRA_YV12 || code == COLOR_YUV2BGR_YV12 || + code == COLOR_YUV2RGBA_YV12 || code == COLOR_YUV2RGB_YV12 ? 1 : 0; + + dstSz = Size(sz.width, sz.height * 2 / 3); + globalsize[0] = dstSz.width / 2; globalsize[1] = (dstSz.height/2 + pxPerWIy - 1) / pxPerWIy; + k.create("YUV2RGB_YV12_IYUV", ocl::imgproc::cvtcolor_oclsrc, + opts + format("-D dcn=%d -D bidx=%d -D uidx=%d%s", dcn, bidx, uidx, + src.isContinuous() ? " -D SRC_CONT" : "")); + break; + } case COLOR_BGR2YCrCb: case COLOR_RGB2YCrCb: { diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index b647f80045..4ac516ae8c 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -301,7 +301,7 @@ __constant int ITUR_BT_601_CVG = 852492; __constant int ITUR_BT_601_CVR = 1673527; __constant int ITUR_BT_601_SHIFT = 20; -__kernel void YUV2RGB_NVx(__global const uchar* srcptr, int src_step, int src_offset, +__kernel void YUV2RGB_NV(__global const uchar* srcptr, int src_step, int src_offset, __global uchar* dstptr, int dst_step, int dt_offset, int rows, int cols) { @@ -318,15 +318,15 @@ __kernel void YUV2RGB_NVx(__global const uchar* srcptr, int src_step, int src_of __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* dst2 = dstptr + mad24((y << 1) + 1, dst_step, x * (dcn<<1) + dt_offset); + __global uchar* dst2 = dst1 + dst_step; int Y1 = ysrc[0]; int Y2 = ysrc[1]; int Y3 = ysrc[src_step]; int Y4 = ysrc[src_step + 1]; - int U = usrc[uidx] - 128; - int V = usrc[1 - uidx] - 128; + int U = ((int)usrc[uidx]) - 128; + int V = ((int)usrc[1-uidx]) - 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; @@ -369,6 +369,83 @@ __kernel void YUV2RGB_NVx(__global const uchar* srcptr, int src_step, int src_of } } +__kernel void YUV2RGB_YV12_IYUV(__global const uchar* srcptr, int src_step, int src_offset, + __global uchar* dstptr, int dst_step, int dt_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) + { + #pragma unroll + for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) + { + if (y < rows / 2 ) + { + __global const uchar* ysrc = srcptr + mad24(y << 1, src_step, (x << 1) + src_offset); + __global uchar* dst1 = dstptr + mad24(y << 1, dst_step, x * (dcn<<1) + dt_offset); + __global uchar* dst2 = dst1 + dst_step; + + int Y1 = ysrc[0]; + int Y2 = ysrc[1]; + int Y3 = ysrc[src_step]; + int Y4 = ysrc[src_step + 1]; + +#ifdef SRC_CONT + __global const uchar* uvsrc = srcptr + mad24(rows, src_step, src_offset); + int u_ind = mad24(y, cols >> 1, x); + int uv[2] = { ((int)uvsrc[u_ind]) - 128, ((int)uvsrc[u_ind + ((rows * cols) >> 2)]) - 128 }; +#else + int vsteps[2] = { cols >> 1, src_step - (cols >> 1)}; + __global const uchar* usrc = srcptr + mad24(rows + (y>>1), src_step, src_offset + (y%2)*(cols >> 1) + x); + __global const uchar* vsrc = usrc + mad24(rows >> 2, src_step, rows % 4 ? vsteps[y%2] : 0); + int uv[2] = { ((int)usrc[0]) - 128, ((int)vsrc[0]) - 128 }; +#endif + int u = uv[uidx]; + int v = uv[1-uidx]; + + 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); +#if dcn == 4 + dst1[3] = 255; +#endif + + Y2 = max(0, Y2 - 16) * ITUR_BT_601_CY; + dst1[dcn + 2 - bidx] = convert_uchar_sat((Y2 + ruv) >> ITUR_BT_601_SHIFT); + dst1[dcn + 1] = convert_uchar_sat((Y2 + guv) >> ITUR_BT_601_SHIFT); + dst1[dcn + bidx] = convert_uchar_sat((Y2 + buv) >> ITUR_BT_601_SHIFT); +#if dcn == 4 + dst1[7] = 255; +#endif + + 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); +#if dcn == 4 + dst2[3] = 255; +#endif + + Y4 = max(0, Y4 - 16) * ITUR_BT_601_CY; + dst2[dcn + 2 - bidx] = convert_uchar_sat((Y4 + ruv) >> ITUR_BT_601_SHIFT); + dst2[dcn + 1] = convert_uchar_sat((Y4 + guv) >> ITUR_BT_601_SHIFT); + dst2[dcn + bidx] = convert_uchar_sat((Y4 + buv) >> ITUR_BT_601_SHIFT); +#if dcn == 4 + dst2[7] = 255; +#endif + } + ++y; + } + } +} + ///////////////////////////////////// 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 2130eca20e..2055326a90 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -320,7 +320,7 @@ OCL_TEST_P(CvtColor8u32f, Luv2RGBA) { performTest(3, 4, CVTCODE(Luv2RGB), depth OCL_TEST_P(CvtColor8u32f, Luv2LBGRA) { performTest(3, 4, CVTCODE(Luv2LBGR), depth == CV_8U ? 1 : 1e-5); } OCL_TEST_P(CvtColor8u32f, Luv2LRGBA) { performTest(3, 4, CVTCODE(Luv2LRGB), depth == CV_8U ? 1 : 1e-5); } -// YUV -> RGBA_NVx +// YUV420 -> RGBA struct CvtColor_YUV420 : public CvtColor @@ -352,6 +352,14 @@ OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV21) { performTest(1, 4, COLOR_YUV2RGBA_NV OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV21) { performTest(1, 4, COLOR_YUV2BGRA_NV21); } OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV21) { performTest(1, 3, COLOR_YUV2RGB_NV21); } OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV21) { performTest(1, 3, COLOR_YUV2BGR_NV21); } +OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_YV12) { performTest(1, 4, COLOR_YUV2RGBA_YV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_YV12) { performTest(1, 4, COLOR_YUV2BGRA_YV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2RGB_YV12) { performTest(1, 3, COLOR_YUV2RGB_YV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2BGR_YV12) { performTest(1, 3, COLOR_YUV2BGR_YV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_IYUV) { performTest(1, 4, COLOR_YUV2RGBA_IYUV); } +OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_IYUV) { performTest(1, 4, COLOR_YUV2BGRA_IYUV); } +OCL_TEST_P(CvtColor_YUV420, YUV2RGB_IYUV) { performTest(1, 3, COLOR_YUV2RGB_IYUV); } +OCL_TEST_P(CvtColor_YUV420, YUV2BGR_IYUV) { performTest(1, 3, COLOR_YUV2BGR_IYUV); } OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u, From 8c91604f5acd60b62e8e7d64f30ff58dc0c582e2 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Fri, 3 Oct 2014 12:29:15 +0400 Subject: [PATCH 06/26] Added OCL code for YUV2GRAY_420 color conversion --- modules/imgproc/src/color.cpp | 14 +++++++++++ modules/imgproc/test/ocl/test_color.cpp | 33 +++++++++++++------------ 2 files changed, 31 insertions(+), 16 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 324b824ae0..a6cebe6bc0 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -4997,6 +4997,20 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) src.isContinuous() ? " -D SRC_CONT" : "")); break; } + case COLOR_YUV2GRAY_420: + { + if (dcn <= 0) dcn = 1; + + CV_Assert( dcn == 1 ); + CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); + + dstSz = Size(sz.width, sz.height * 2 / 3); + _dst.create(dstSz, CV_MAKETYPE(depth, dcn)); + dst = _dst.getUMat(); + + src.rowRange(0, dstSz.height).copyTo(dst); + return true; + } case COLOR_BGR2YCrCb: case COLOR_RGB2YCrCb: { diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index 2055326a90..43343469c5 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -344,22 +344,23 @@ struct CvtColor_YUV420 : } }; -OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) { performTest(1, 4, COLOR_YUV2RGBA_NV12); } -OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) { performTest(1, 4, COLOR_YUV2BGRA_NV12); } -OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { performTest(1, 3, COLOR_YUV2RGB_NV12); } -OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { performTest(1, 3, COLOR_YUV2BGR_NV12); } -OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV21) { performTest(1, 4, COLOR_YUV2RGBA_NV21); } -OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV21) { performTest(1, 4, COLOR_YUV2BGRA_NV21); } -OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV21) { performTest(1, 3, COLOR_YUV2RGB_NV21); } -OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV21) { performTest(1, 3, COLOR_YUV2BGR_NV21); } -OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_YV12) { performTest(1, 4, COLOR_YUV2RGBA_YV12); } -OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_YV12) { performTest(1, 4, COLOR_YUV2BGRA_YV12); } -OCL_TEST_P(CvtColor_YUV420, YUV2RGB_YV12) { performTest(1, 3, COLOR_YUV2RGB_YV12); } -OCL_TEST_P(CvtColor_YUV420, YUV2BGR_YV12) { performTest(1, 3, COLOR_YUV2BGR_YV12); } -OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_IYUV) { performTest(1, 4, COLOR_YUV2RGBA_IYUV); } -OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_IYUV) { performTest(1, 4, COLOR_YUV2BGRA_IYUV); } -OCL_TEST_P(CvtColor_YUV420, YUV2RGB_IYUV) { performTest(1, 3, COLOR_YUV2RGB_IYUV); } -OCL_TEST_P(CvtColor_YUV420, YUV2BGR_IYUV) { performTest(1, 3, COLOR_YUV2BGR_IYUV); } +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_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u, From c8707b891b1cd7ca9f717ce95598cdab26865aa3 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Mon, 6 Oct 2014 19:19:44 +0400 Subject: [PATCH 07/26] 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())); From 5aa9ac9a7788de69726bf696cc6e3ffd14f7d895 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Mon, 6 Oct 2014 19:21:57 +0400 Subject: [PATCH 08/26] Added OCL code for YUV422 -> RGB[A]|BGR[A] color conversion --- modules/imgproc/src/color.cpp | 32 ++++++-- modules/imgproc/src/opencl/cvtcolor.cl | 105 +++++++++++++++++------- modules/imgproc/test/ocl/test_color.cpp | 43 ++++++++++ 3 files changed, 146 insertions(+), 34 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 6cca73652f..f363189579 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -4848,7 +4848,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) bool ok = false; UMat src = _src.getUMat(), dst; Size sz = src.size(), dstSz = sz; - int scn = src.channels(), depth = src.depth(), bidx, uidx; + int scn = src.channels(), depth = src.depth(), bidx, uidx, yidx; int dims = 2, stripeSize = 1; ocl::Kernel k; @@ -4967,14 +4967,14 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); dcn = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2RGBA_NV12 || code == COLOR_YUV2BGRA_NV21 || code == COLOR_YUV2RGBA_NV21 ? 4 : 3; - bidx = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2BGR_NV12 || + bidx = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2BGR_NV12 || code == COLOR_YUV2BGRA_NV21 || code == COLOR_YUV2BGR_NV21 ? 0 : 2; uidx = code == COLOR_YUV2RGBA_NV21 || code == COLOR_YUV2RGB_NV21 || code == COLOR_YUV2BGRA_NV21 || code == COLOR_YUV2BGR_NV21 ? 1 : 0; dstSz = Size(sz.width, sz.height * 2 / 3); globalsize[0] = dstSz.width / 2; globalsize[1] = (dstSz.height/2 + pxPerWIy - 1) / pxPerWIy; - k.create("YUV2RGB_NV", ocl::imgproc::cvtcolor_oclsrc, + k.create("YUV2RGB_NVx", ocl::imgproc::cvtcolor_oclsrc, opts + format("-D dcn=%d -D bidx=%d -D uidx=%d", dcn, bidx, uidx)); break; } @@ -4985,7 +4985,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); dcn = code == COLOR_YUV2BGRA_YV12 || code == COLOR_YUV2RGBA_YV12 || code == COLOR_YUV2BGRA_IYUV || code == COLOR_YUV2RGBA_IYUV ? 4 : 3; - bidx = code == COLOR_YUV2BGRA_YV12 || code == COLOR_YUV2BGR_YV12 || + bidx = code == COLOR_YUV2BGRA_YV12 || code == COLOR_YUV2BGR_YV12 || code == COLOR_YUV2BGRA_IYUV || code == COLOR_YUV2BGR_IYUV ? 0 : 2; uidx = code == COLOR_YUV2BGRA_YV12 || code == COLOR_YUV2BGR_YV12 || code == COLOR_YUV2RGBA_YV12 || code == COLOR_YUV2RGB_YV12 ? 1 : 0; @@ -5015,7 +5015,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) 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 || + 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; @@ -5030,6 +5030,28 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) opts + format("-D dcn=%d -D bidx=%d -D uidx=%d", dcn, bidx, uidx)); break; } + case COLOR_YUV2RGB_UYVY: case COLOR_YUV2BGR_UYVY: case COLOR_YUV2RGBA_UYVY: case COLOR_YUV2BGRA_UYVY: + case COLOR_YUV2RGB_YUY2: case COLOR_YUV2BGR_YUY2: case COLOR_YUV2RGB_YVYU: case COLOR_YUV2BGR_YVYU: + case COLOR_YUV2RGBA_YUY2: case COLOR_YUV2BGRA_YUY2: case COLOR_YUV2RGBA_YVYU: case COLOR_YUV2BGRA_YVYU: + { + if (dcn <= 0) + dcn = (code==COLOR_YUV2RGBA_UYVY || code==COLOR_YUV2BGRA_UYVY || code==COLOR_YUV2RGBA_YUY2 || + code==COLOR_YUV2BGRA_YUY2 || code==COLOR_YUV2RGBA_YVYU || code==COLOR_YUV2BGRA_YVYU) ? 4 : 3; + + bidx = (code==COLOR_YUV2BGR_UYVY || code==COLOR_YUV2BGRA_UYVY || code==COLOR_YUV2BGRA_YUY2 || + code==COLOR_YUV2BGR_YUY2 || code==COLOR_YUV2BGRA_YVYU || code==COLOR_YUV2BGR_YVYU) ? 0 : 2; + yidx = (code==COLOR_YUV2RGB_UYVY || code==COLOR_YUV2RGBA_UYVY || code==COLOR_YUV2BGR_UYVY || code==COLOR_YUV2BGRA_UYVY) ? 1 : 0; + uidx = (code==COLOR_YUV2RGB_YVYU || code==COLOR_YUV2RGBA_YVYU || + code==COLOR_YUV2BGR_YVYU || code==COLOR_YUV2BGRA_YVYU) ? 2 : 0; + uidx = 1 - yidx + uidx; + + CV_Assert( dcn == 3 || dcn == 4 ); + CV_Assert( scn == 2 && depth == CV_8U ); + + k.create("YUV2RGB_422", ocl::imgproc::cvtcolor_oclsrc, + opts + format("-D dcn=%d -D bidx=%d -D uidx=%d -D yidx=%d", dcn, bidx, uidx, yidx)); + break; + } case COLOR_BGR2YCrCb: case COLOR_RGB2YCrCb: { diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index f57c0c068c..cf7c06ee68 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -300,12 +300,12 @@ __kernel void YUV2RGB(__global const uchar* srcptr, int src_step, int src_offset __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_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 YUV2RGB_NV(__global const uchar* srcptr, int src_step, int src_offset, +__kernel void YUV2RGB_NVx(__global const uchar* srcptr, int src_step, int src_offset, __global uchar* dstptr, int dst_step, int dt_offset, int rows, int cols) { @@ -329,41 +329,41 @@ __kernel void YUV2RGB_NV(__global const uchar* srcptr, int src_step, int src_off int Y3 = ysrc[src_step]; int Y4 = ysrc[src_step + 1]; - int U = ((int)usrc[uidx]) - 128; - int V = ((int)usrc[1-uidx]) - 128; + int U = ((int)usrc[uidx]) - HALF_MAX; + int V = ((int)usrc[1-uidx]) - HALF_MAX; - 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; + int ruv = mad24(ITUR_BT_601_CVR, V, (1 << (ITUR_BT_601_SHIFT - 1))); + int guv = mad24(ITUR_BT_601_CVG, V, mad24(ITUR_BT_601_CUG, U, (1 << (ITUR_BT_601_SHIFT - 1)))); + int buv = mad24(ITUR_BT_601_CUB, U, (1 << (ITUR_BT_601_SHIFT - 1))); - Y1 = max(0, Y1 - 16) * ITUR_BT_601_CY; - dst1[2 - bidx] = convert_uchar_sat((Y1 + ruv) >> ITUR_BT_601_SHIFT); + Y1 = mul24(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[bidx] = convert_uchar_sat((Y1 + buv) >> ITUR_BT_601_SHIFT); #if dcn == 4 dst1[3] = 255; #endif - Y2 = max(0, Y2 - 16) * ITUR_BT_601_CY; + Y2 = mul24(max(0, Y2 - 16), ITUR_BT_601_CY); dst1[dcn + 2 - bidx] = convert_uchar_sat((Y2 + ruv) >> ITUR_BT_601_SHIFT); dst1[dcn + 1] = convert_uchar_sat((Y2 + guv) >> ITUR_BT_601_SHIFT); - dst1[dcn + bidx] = convert_uchar_sat((Y2 + buv) >> ITUR_BT_601_SHIFT); + dst1[dcn + bidx] = convert_uchar_sat((Y2 + buv) >> ITUR_BT_601_SHIFT); #if dcn == 4 dst1[7] = 255; #endif - Y3 = max(0, Y3 - 16) * ITUR_BT_601_CY; - dst2[2 - bidx] = convert_uchar_sat((Y3 + ruv) >> ITUR_BT_601_SHIFT); + Y3 = mul24(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[bidx] = convert_uchar_sat((Y3 + buv) >> ITUR_BT_601_SHIFT); #if dcn == 4 dst2[3] = 255; #endif - Y4 = max(0, Y4 - 16) * ITUR_BT_601_CY; + Y4 = mul24(max(0, Y4 - 16), ITUR_BT_601_CY); dst2[dcn + 2 - bidx] = convert_uchar_sat((Y4 + ruv) >> ITUR_BT_601_SHIFT); dst2[dcn + 1] = convert_uchar_sat((Y4 + guv) >> ITUR_BT_601_SHIFT); - dst2[dcn + bidx] = convert_uchar_sat((Y4 + buv) >> ITUR_BT_601_SHIFT); + dst2[dcn + bidx] = convert_uchar_sat((Y4 + buv) >> ITUR_BT_601_SHIFT); #if dcn == 4 dst2[7] = 255; #endif @@ -399,21 +399,21 @@ __kernel void YUV2RGB_YV12_IYUV(__global const uchar* srcptr, int src_step, int #ifdef SRC_CONT __global const uchar* uvsrc = srcptr + mad24(rows, src_step, src_offset); int u_ind = mad24(y, cols >> 1, x); - int uv[2] = { ((int)uvsrc[u_ind]) - 128, ((int)uvsrc[u_ind + ((rows * cols) >> 2)]) - 128 }; + int uv[2] = { ((int)uvsrc[u_ind]) - HALF_MAX, ((int)uvsrc[u_ind + ((rows * cols) >> 2)]) - HALF_MAX }; #else int vsteps[2] = { cols >> 1, src_step - (cols >> 1)}; __global const uchar* usrc = srcptr + mad24(rows + (y>>1), src_step, src_offset + (y%2)*(cols >> 1) + x); __global const uchar* vsrc = usrc + mad24(rows >> 2, src_step, rows % 4 ? vsteps[y%2] : 0); - int uv[2] = { ((int)usrc[0]) - 128, ((int)vsrc[0]) - 128 }; + int uv[2] = { ((int)usrc[0]) - HALF_MAX, ((int)vsrc[0]) - HALF_MAX }; #endif - int u = uv[uidx]; - int v = uv[1-uidx]; + int U = uv[uidx]; + int V = uv[1-uidx]; - 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; + int ruv = mad24(ITUR_BT_601_CVR, V, (1 << (ITUR_BT_601_SHIFT - 1))); + int guv = mad24(ITUR_BT_601_CVG, V, mad24(ITUR_BT_601_CUG, U, (1 << (ITUR_BT_601_SHIFT - 1)))); + int buv = mad24(ITUR_BT_601_CUB, U, (1 << (ITUR_BT_601_SHIFT - 1))); - Y1 = max(0, Y1 - 16) * ITUR_BT_601_CY; + Y1 = mul24(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); @@ -421,7 +421,7 @@ __kernel void YUV2RGB_YV12_IYUV(__global const uchar* srcptr, int src_step, int dst1[3] = 255; #endif - Y2 = max(0, Y2 - 16) * ITUR_BT_601_CY; + Y2 = mul24(max(0, Y2 - 16), ITUR_BT_601_CY); dst1[dcn + 2 - bidx] = convert_uchar_sat((Y2 + ruv) >> ITUR_BT_601_SHIFT); dst1[dcn + 1] = convert_uchar_sat((Y2 + guv) >> ITUR_BT_601_SHIFT); dst1[dcn + bidx] = convert_uchar_sat((Y2 + buv) >> ITUR_BT_601_SHIFT); @@ -429,7 +429,7 @@ __kernel void YUV2RGB_YV12_IYUV(__global const uchar* srcptr, int src_step, int dst1[7] = 255; #endif - Y3 = max(0, Y3 - 16) * ITUR_BT_601_CY; + Y3 = mul24(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); @@ -437,7 +437,7 @@ __kernel void YUV2RGB_YV12_IYUV(__global const uchar* srcptr, int src_step, int dst2[3] = 255; #endif - Y4 = max(0, Y4 - 16) * ITUR_BT_601_CY; + Y4 = mul24(max(0, Y4 - 16), ITUR_BT_601_CY); dst2[dcn + 2 - bidx] = convert_uchar_sat((Y4 + ruv) >> ITUR_BT_601_SHIFT); dst2[dcn + 1] = convert_uchar_sat((Y4 + guv) >> ITUR_BT_601_SHIFT); dst2[dcn + bidx] = convert_uchar_sat((Y4 + buv) >> ITUR_BT_601_SHIFT); @@ -517,6 +517,53 @@ __kernel void RGB2YUV_YV12_IYUV(__global const uchar* srcptr, int src_step, int } } +__kernel void YUV2RGB_422(__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) + { + __global const uchar* src = srcptr + mad24(y, src_step, (x << 2) + src_offset); + __global uchar* dst = dstptr + mad24(y, dst_step, mad24(x << 1, dcn, dst_offset)); + + #pragma unroll + for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) + { + if (y < rows ) + { + int U = ((int) src[uidx]) - HALF_MAX; + int V = ((int) src[(2 + uidx) % 4]) - HALF_MAX; + + int ruv = mad24(ITUR_BT_601_CVR, V, (1 << (ITUR_BT_601_SHIFT - 1))); + int guv = mad24(ITUR_BT_601_CVG, V, mad24(ITUR_BT_601_CUG, U, (1 << (ITUR_BT_601_SHIFT - 1)))); + int buv = mad24(ITUR_BT_601_CUB, U, (1 << (ITUR_BT_601_SHIFT - 1))); + + int y00 = max(0, ((int) src[yidx]) - 16) * ITUR_BT_601_CY; + dst[2 - bidx] = convert_uchar_sat((y00 + ruv) >> ITUR_BT_601_SHIFT); + dst[1] = convert_uchar_sat((y00 + guv) >> ITUR_BT_601_SHIFT); + dst[bidx] = convert_uchar_sat((y00 + buv) >> ITUR_BT_601_SHIFT); +#if dcn == 4 + dst[3] = 255; +#endif + + int y01 = max(0, ((int) src[yidx + 2]) - 16) * ITUR_BT_601_CY; + dst[dcn + 2 - bidx] = convert_uchar_sat((y01 + ruv) >> ITUR_BT_601_SHIFT); + dst[dcn + 1] = convert_uchar_sat((y01 + guv) >> ITUR_BT_601_SHIFT); + dst[dcn + bidx] = convert_uchar_sat((y01 + buv) >> ITUR_BT_601_SHIFT); +#if dcn == 4 + dst[7] = 255; +#endif + } + ++y; + src += src_step; + dst += 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 1484db9720..89affcfac7 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -396,6 +396,44 @@ OCL_TEST_P(CvtColor_RGB2YUV_420, BGRA2YUV_IYUV) { performTest(4, 1, CVTCODE(BGRA 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)); } +// YUV422 -> RGBA + +struct CvtColor_YUV2RGB_422 : + 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 roiSize = randomSize(1, MAX_VALUE); + roiSize.width *= 2; + + Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src, src_roi, roiSize, srcBorder, srcType, 2, 100); + + Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, dstType, 5, 16); + + UMAT_UPLOAD_INPUT_PARAMETER(src); + UMAT_UPLOAD_OUTPUT_PARAMETER(dst); + } +}; + +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2RGB_UYVY) { performTest(2, 3, CVTCODE(YUV2RGB_UYVY)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2BGR_UYVY) { performTest(2, 3, CVTCODE(YUV2BGR_UYVY)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2RGBA_UYVY) { performTest(2, 4, CVTCODE(YUV2RGBA_UYVY)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2BGRA_UYVY) { performTest(2, 4, CVTCODE(YUV2BGRA_UYVY)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2RGB_YUY2) { performTest(2, 3, CVTCODE(YUV2RGB_YUY2)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2BGR_YUY2) { performTest(2, 3, CVTCODE(YUV2BGR_YUY2)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2RGBA_YUY2) { performTest(2, 4, CVTCODE(YUV2RGBA_YUY2)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2BGRA_YUY2) { performTest(2, 4, CVTCODE(YUV2BGRA_YUY2)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2RGB_YVYU) { performTest(2, 3, CVTCODE(YUV2RGB_YVYU)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2BGR_YVYU) { performTest(2, 3, CVTCODE(YUV2BGR_YVYU)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2RGBA_YVYU) { performTest(2, 4, CVTCODE(YUV2RGBA_YVYU)); } +OCL_TEST_P(CvtColor_YUV2RGB_422, YUV2BGRA_YVYU) { performTest(2, 4, CVTCODE(YUV2BGRA_YVYU)); } + + OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u, testing::Combine(testing::Values(MatDepth(CV_8U)), Bool())); @@ -417,6 +455,11 @@ OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor_RGB2YUV_420, testing::Values(MatDepth(CV_8U)), Bool())); +OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor_YUV2RGB_422, + testing::Combine( + testing::Values(MatDepth(CV_8U)), + Bool())); + } } // namespace cvtest::ocl #endif From 60367907fe75d369dc3158b4a696371bfa2ce6ee Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Tue, 14 Oct 2014 16:31:10 +0400 Subject: [PATCH 09/26] Used direct float calculations --- modules/imgproc/src/opencl/cvtcolor.cl | 205 +++++++++++------------- modules/imgproc/test/ocl/test_color.cpp | 16 +- 2 files changed, 104 insertions(+), 117 deletions(-) diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index cf7c06ee68..c3cfd0d592 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -77,7 +77,7 @@ enum { yuv_shift = 14, xyz_shift = 12, - hsv_shift = 12, + hsv_shift = 12, R2Y = 4899, G2Y = 9617, B2Y = 1868, @@ -149,7 +149,7 @@ __kernel void RGB2Gray(__global const uchar * srcptr, int src_step, int src_offs #ifdef DEPTH_5 dst[0] = fma(src_pix.B_COMP, 0.114f, fma(src_pix.G_COMP, 0.587f, src_pix.R_COMP * 0.299f)); #else - dst[0] = (DATA_TYPE)CV_DESCALE(mad24(src_pix.B_COMP, B2Y, mad24(src_pix.G_COMP, G2Y, src_pix.R_COMP * R2Y)), yuv_shift); + dst[0] = (DATA_TYPE)CV_DESCALE(mad24(src_pix.B_COMP, B2Y, mad24(src_pix.G_COMP, G2Y, mul24(src_pix.R_COMP, R2Y))), yuv_shift); #endif ++y; src_index += src_step; @@ -224,13 +224,13 @@ __kernel void RGB2YUV(__global const uchar* srcptr, int src_step, int src_offset #ifdef DEPTH_5 __constant float * coeffs = c_RGB2YUVCoeffs_f; - const DATA_TYPE Y = fma(b, coeffs[0], fma(g, coeffs[1], r * coeffs[2])); + const DATA_TYPE Y = fma(b, coeffs[0], fma(g, coeffs[1], r * coeffs[2])); const DATA_TYPE U = fma(b - Y, coeffs[3], HALF_MAX); const DATA_TYPE V = fma(r - 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(mad24(b, coeffs[0], mad24(g, coeffs[1], r * coeffs[2])), yuv_shift); + const int Y = CV_DESCALE(mad24(b, coeffs[0], mad24(g, coeffs[1], mul24(r, coeffs[2]))), yuv_shift); const int U = CV_DESCALE(mad24(b - Y, coeffs[3], delta), yuv_shift); const int V = CV_DESCALE(mad24(r - Y, coeffs[4], delta), yuv_shift); #endif @@ -247,8 +247,8 @@ __kernel void RGB2YUV(__global const uchar* srcptr, int src_step, int src_offset } } -__constant float c_YUV2RGBCoeffs_f[5] = { 2.032f, -0.395f, -0.581f, 1.140f }; -__constant int c_YUV2RGBCoeffs_i[5] = { 33292, -6472, -9519, 18678 }; +__constant float c_YUV2RGBCoeffs_f[4] = { 2.032f, -0.395f, -0.581f, 1.140f }; +__constant int c_YUV2RGBCoeffs_i[4] = { 33292, -6472, -9519, 18678 }; __kernel void YUV2RGB(__global const uchar* srcptr, int src_step, int src_offset, __global uchar* dstptr, int dst_step, int dt_offset, @@ -279,9 +279,9 @@ __kernel void YUV2RGB(__global const uchar* srcptr, int src_step, int src_offset float b = fma(U - HALF_MAX, coeffs[0], Y); #else __constant int * coeffs = c_YUV2RGBCoeffs_i; - const int r = Y + CV_DESCALE((V - HALF_MAX) * coeffs[3], yuv_shift); - const int g = Y + CV_DESCALE(mad24(V - HALF_MAX, coeffs[2], (U - HALF_MAX) * coeffs[1]), yuv_shift); - const int b = Y + CV_DESCALE((U - HALF_MAX) * coeffs[0], yuv_shift); + const int r = Y + CV_DESCALE(mul24(V - HALF_MAX, coeffs[3]), yuv_shift); + const int g = Y + CV_DESCALE(mad24(V - HALF_MAX, coeffs[2], mul24(U - HALF_MAX, coeffs[1])), yuv_shift); + const int b = Y + CV_DESCALE(mul24(U - HALF_MAX, coeffs[0]), yuv_shift); #endif dst[bidx] = SAT_CAST( b ); @@ -297,13 +297,8 @@ __kernel void YUV2RGB(__global const uchar* srcptr, int src_step, int src_offset } } } - -__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; +__constant float c_YUV2RGBCoeffs_420[5] = { 1.163999557f, 2.017999649f, -0.390999794f, + -0.812999725f, 1.5959997177f }; __kernel void YUV2RGB_NVx(__global const uchar* srcptr, int src_step, int src_offset, __global uchar* dstptr, int dst_step, int dt_offset, @@ -324,46 +319,47 @@ __kernel void YUV2RGB_NVx(__global const uchar* srcptr, int src_step, int src_of __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]; - int Y2 = ysrc[1]; - int Y3 = ysrc[src_step]; - int Y4 = ysrc[src_step + 1]; + float Y1 = ysrc[0]; + float Y2 = ysrc[1]; + float Y3 = ysrc[src_step]; + float Y4 = ysrc[src_step + 1]; - int U = ((int)usrc[uidx]) - HALF_MAX; - int V = ((int)usrc[1-uidx]) - HALF_MAX; + float U = ((float)usrc[uidx]) - HALF_MAX; + float V = ((float)usrc[1-uidx]) - HALF_MAX; - int ruv = mad24(ITUR_BT_601_CVR, V, (1 << (ITUR_BT_601_SHIFT - 1))); - int guv = mad24(ITUR_BT_601_CVG, V, mad24(ITUR_BT_601_CUG, U, (1 << (ITUR_BT_601_SHIFT - 1)))); - int buv = mad24(ITUR_BT_601_CUB, U, (1 << (ITUR_BT_601_SHIFT - 1))); + __constant float* coeffs = c_YUV2RGBCoeffs_420; + float ruv = fma(coeffs[4], V, 0.5f); + float guv = fma(coeffs[3], V, fma(coeffs[2], U, 0.5f)); + float buv = fma(coeffs[1], U, 0.5f); - Y1 = mul24(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); + Y1 = max(0.f, Y1 - 16.f) * coeffs[0]; + dst1[2 - bidx] = convert_uchar_sat(Y1 + ruv); + dst1[1] = convert_uchar_sat(Y1 + guv); + dst1[bidx] = convert_uchar_sat(Y1 + buv); #if dcn == 4 dst1[3] = 255; #endif - Y2 = mul24(max(0, Y2 - 16), ITUR_BT_601_CY); - dst1[dcn + 2 - bidx] = convert_uchar_sat((Y2 + ruv) >> ITUR_BT_601_SHIFT); - dst1[dcn + 1] = convert_uchar_sat((Y2 + guv) >> ITUR_BT_601_SHIFT); - dst1[dcn + bidx] = convert_uchar_sat((Y2 + buv) >> ITUR_BT_601_SHIFT); + Y2 = max(0.f, Y2 - 16.f) * coeffs[0]; + dst1[dcn + 2 - bidx] = convert_uchar_sat(Y2 + ruv); + dst1[dcn + 1] = convert_uchar_sat(Y2 + guv); + dst1[dcn + bidx] = convert_uchar_sat(Y2 + buv); #if dcn == 4 dst1[7] = 255; #endif - Y3 = mul24(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); + Y3 = max(0.f, Y3 - 16.f) * coeffs[0]; + dst2[2 - bidx] = convert_uchar_sat(Y3 + ruv); + dst2[1] = convert_uchar_sat(Y3 + guv); + dst2[bidx] = convert_uchar_sat(Y3 + buv); #if dcn == 4 dst2[3] = 255; #endif - Y4 = mul24(max(0, Y4 - 16), ITUR_BT_601_CY); - dst2[dcn + 2 - bidx] = convert_uchar_sat((Y4 + ruv) >> ITUR_BT_601_SHIFT); - dst2[dcn + 1] = convert_uchar_sat((Y4 + guv) >> ITUR_BT_601_SHIFT); - dst2[dcn + bidx] = convert_uchar_sat((Y4 + buv) >> ITUR_BT_601_SHIFT); + Y4 = max(0.f, Y4 - 16.f) * coeffs[0]; + dst2[dcn + 2 - bidx] = convert_uchar_sat(Y4 + ruv); + dst2[dcn + 1] = convert_uchar_sat(Y4 + guv); + dst2[dcn + bidx] = convert_uchar_sat(Y4 + buv); #if dcn == 4 dst2[7] = 255; #endif @@ -391,56 +387,57 @@ __kernel void YUV2RGB_YV12_IYUV(__global const uchar* srcptr, int src_step, int __global uchar* dst1 = dstptr + mad24(y << 1, dst_step, x * (dcn<<1) + dt_offset); __global uchar* dst2 = dst1 + dst_step; - int Y1 = ysrc[0]; - int Y2 = ysrc[1]; - int Y3 = ysrc[src_step]; - int Y4 = ysrc[src_step + 1]; + float Y1 = ysrc[0]; + float Y2 = ysrc[1]; + float Y3 = ysrc[src_step]; + float Y4 = ysrc[src_step + 1]; #ifdef SRC_CONT __global const uchar* uvsrc = srcptr + mad24(rows, src_step, src_offset); int u_ind = mad24(y, cols >> 1, x); - int uv[2] = { ((int)uvsrc[u_ind]) - HALF_MAX, ((int)uvsrc[u_ind + ((rows * cols) >> 2)]) - HALF_MAX }; + float uv[2] = { ((float)uvsrc[u_ind]) - HALF_MAX, ((float)uvsrc[u_ind + ((rows * cols) >> 2)]) - HALF_MAX }; #else int vsteps[2] = { cols >> 1, src_step - (cols >> 1)}; __global const uchar* usrc = srcptr + mad24(rows + (y>>1), src_step, src_offset + (y%2)*(cols >> 1) + x); __global const uchar* vsrc = usrc + mad24(rows >> 2, src_step, rows % 4 ? vsteps[y%2] : 0); - int uv[2] = { ((int)usrc[0]) - HALF_MAX, ((int)vsrc[0]) - HALF_MAX }; + float uv[2] = { ((float)usrc[0]) - HALF_MAX, ((float)vsrc[0]) - HALF_MAX }; #endif - int U = uv[uidx]; - int V = uv[1-uidx]; + float U = uv[uidx]; + float V = uv[1-uidx]; - int ruv = mad24(ITUR_BT_601_CVR, V, (1 << (ITUR_BT_601_SHIFT - 1))); - int guv = mad24(ITUR_BT_601_CVG, V, mad24(ITUR_BT_601_CUG, U, (1 << (ITUR_BT_601_SHIFT - 1)))); - int buv = mad24(ITUR_BT_601_CUB, U, (1 << (ITUR_BT_601_SHIFT - 1))); + __constant float* coeffs = c_YUV2RGBCoeffs_420; + float ruv = fma(coeffs[4], V, 0.5f); + float guv = fma(coeffs[3], V, fma(coeffs[2], U, 0.5f)); + float buv = fma(coeffs[1], U, 0.5f); - Y1 = mul24(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); + Y1 = max(0.f, Y1 - 16.f) * coeffs[0]; + dst1[2 - bidx] = convert_uchar_sat(Y1 + ruv); + dst1[1] = convert_uchar_sat(Y1 + guv); + dst1[bidx] = convert_uchar_sat(Y1 + buv); #if dcn == 4 dst1[3] = 255; #endif - Y2 = mul24(max(0, Y2 - 16), ITUR_BT_601_CY); - dst1[dcn + 2 - bidx] = convert_uchar_sat((Y2 + ruv) >> ITUR_BT_601_SHIFT); - dst1[dcn + 1] = convert_uchar_sat((Y2 + guv) >> ITUR_BT_601_SHIFT); - dst1[dcn + bidx] = convert_uchar_sat((Y2 + buv) >> ITUR_BT_601_SHIFT); + Y2 = max(0.f, Y2 - 16.f) * coeffs[0]; + dst1[dcn + 2 - bidx] = convert_uchar_sat(Y2 + ruv); + dst1[dcn + 1] = convert_uchar_sat(Y2 + guv); + dst1[dcn + bidx] = convert_uchar_sat(Y2 + buv); #if dcn == 4 dst1[7] = 255; #endif - Y3 = mul24(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); + Y3 = max(0.f, Y3 - 16.f) * coeffs[0]; + dst2[2 - bidx] = convert_uchar_sat(Y3 + ruv); + dst2[1] = convert_uchar_sat(Y3 + guv); + dst2[bidx] = convert_uchar_sat(Y3 + buv); #if dcn == 4 dst2[3] = 255; #endif - Y4 = mul24(max(0, Y4 - 16), ITUR_BT_601_CY); - dst2[dcn + 2 - bidx] = convert_uchar_sat((Y4 + ruv) >> ITUR_BT_601_SHIFT); - dst2[dcn + 1] = convert_uchar_sat((Y4 + guv) >> ITUR_BT_601_SHIFT); - dst2[dcn + bidx] = convert_uchar_sat((Y4 + buv) >> ITUR_BT_601_SHIFT); + Y4 = max(0.f, Y4 - 16.f) * coeffs[0]; + dst2[dcn + 2 - bidx] = convert_uchar_sat(Y4 + ruv); + dst2[dcn + 1] = convert_uchar_sat(Y4 + guv); + dst2[dcn + bidx] = convert_uchar_sat(Y4 + buv); #if dcn == 4 dst2[7] = 255; #endif @@ -450,16 +447,8 @@ __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; +__constant float c_RGB2YUVCoeffs_420[8] = { 0.256999969f, 0.50399971f, 0.09799957f, -0.1479988098f, -0.2909994125f, + 0.438999176f, -0.3679990768f, -0.0709991455f }; __kernel void RGB2YUV_YV12_IYUV(__global const uchar* srcptr, int src_step, int src_offset, __global uchar* dstptr, int dst_step, int dst_offset, @@ -488,26 +477,22 @@ __kernel void RGB2YUV_YV12_IYUV(__global const uchar* srcptr, int src_step, int __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)); + float4 src_pix1 = convert_float4(vload4(0, src1)); + float4 src_pix2 = convert_float4(vload4(0, src1+scn)); + float4 src_pix3 = convert_float4(vload4(0, src2)); + float4 src_pix4 = convert_float4(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))); + __constant float* coeffs = c_RGB2YUVCoeffs_420; + ydst1[0] = convert_uchar_sat(fma(coeffs[0], src_pix1.R_COMP, fma(coeffs[1], src_pix1.G_COMP, fma(coeffs[2], src_pix1.B_COMP, 16.5f)))); + ydst1[1] = convert_uchar_sat(fma(coeffs[0], src_pix2.R_COMP, fma(coeffs[1], src_pix2.G_COMP, fma(coeffs[2], src_pix2.B_COMP, 16.5f)))); + ydst2[0] = convert_uchar_sat(fma(coeffs[0], src_pix3.R_COMP, fma(coeffs[1], src_pix3.G_COMP, fma(coeffs[2], src_pix3.B_COMP, 16.5f)))); + ydst2[1] = convert_uchar_sat(fma(coeffs[0], src_pix4.R_COMP, fma(coeffs[1], src_pix4.G_COMP, fma(coeffs[2], src_pix4.B_COMP, 16.5f)))); - 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); + float uv[2] = { fma(coeffs[3], src_pix1.R_COMP, fma(coeffs[4], src_pix1.G_COMP, fma(coeffs[5], src_pix1.B_COMP, 128.5f))), + fma(coeffs[5], src_pix1.R_COMP, fma(coeffs[6], src_pix1.G_COMP, fma(coeffs[7], src_pix1.B_COMP, 128.5f))) }; - 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); + udst[0] = convert_uchar_sat(uv[uidx] ); + vdst[0] = convert_uchar_sat(uv[1-uidx]); ++y; src_index += 2*src_step; @@ -534,25 +519,27 @@ __kernel void YUV2RGB_422(__global const uchar* srcptr, int src_step, int src_of { if (y < rows ) { - int U = ((int) src[uidx]) - HALF_MAX; - int V = ((int) src[(2 + uidx) % 4]) - HALF_MAX; + float U = ((float) src[uidx]) - HALF_MAX; + float V = ((float) src[(2 + uidx) % 4]) - HALF_MAX; - int ruv = mad24(ITUR_BT_601_CVR, V, (1 << (ITUR_BT_601_SHIFT - 1))); - int guv = mad24(ITUR_BT_601_CVG, V, mad24(ITUR_BT_601_CUG, U, (1 << (ITUR_BT_601_SHIFT - 1)))); - int buv = mad24(ITUR_BT_601_CUB, U, (1 << (ITUR_BT_601_SHIFT - 1))); - int y00 = max(0, ((int) src[yidx]) - 16) * ITUR_BT_601_CY; - dst[2 - bidx] = convert_uchar_sat((y00 + ruv) >> ITUR_BT_601_SHIFT); - dst[1] = convert_uchar_sat((y00 + guv) >> ITUR_BT_601_SHIFT); - dst[bidx] = convert_uchar_sat((y00 + buv) >> ITUR_BT_601_SHIFT); + __constant float* coeffs = c_YUV2RGBCoeffs_420; + float ruv = fma(coeffs[4], V, 0.5f); + float guv = fma(coeffs[3], V, fma(coeffs[2], U, 0.5f)); + float buv = fma(coeffs[1], U, 0.5f); + + float y00 = max(0.f, ((float) src[yidx]) - 16.f) * coeffs[0]; + dst[2 - bidx] = convert_uchar_sat(y00 + ruv); + dst[1] = convert_uchar_sat(y00 + guv); + dst[bidx] = convert_uchar_sat(y00 + buv); #if dcn == 4 dst[3] = 255; #endif - int y01 = max(0, ((int) src[yidx + 2]) - 16) * ITUR_BT_601_CY; - dst[dcn + 2 - bidx] = convert_uchar_sat((y01 + ruv) >> ITUR_BT_601_SHIFT); - dst[dcn + 1] = convert_uchar_sat((y01 + guv) >> ITUR_BT_601_SHIFT); - dst[dcn + bidx] = convert_uchar_sat((y01 + buv) >> ITUR_BT_601_SHIFT); + float y01 = max(0.f, ((float) src[yidx + 2]) - 16.f) * coeffs[0]; + dst[dcn + 2 - bidx] = convert_uchar_sat(y01 + ruv); + dst[dcn + 1] = convert_uchar_sat(y01 + guv); + dst[dcn + bidx] = convert_uchar_sat(y01 + buv); #if dcn == 4 dst[7] = 255; #endif @@ -599,7 +586,7 @@ __kernel void RGB2YCrCb(__global const uchar* srcptr, int src_step, int src_offs #else __constant int * coeffs = c_RGB2YCrCbCoeffs_i; int delta = HALF_MAX * (1 << yuv_shift); - int Y = CV_DESCALE(mad24(b, coeffs[2], mad24(g, coeffs[1], r * coeffs[0])), yuv_shift); + int Y = CV_DESCALE(mad24(b, coeffs[2], mad24(g, coeffs[1], mul24(r, coeffs[0]))), yuv_shift); int Cr = CV_DESCALE(mad24(r - Y, coeffs[3], delta), yuv_shift); int Cb = CV_DESCALE(mad24(b - Y, coeffs[4], delta), yuv_shift); #endif diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index 89affcfac7..c53607fdbd 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -387,14 +387,14 @@ struct CvtColor_RGB2YUV_420 : } }; -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_TEST_P(CvtColor_RGB2YUV_420, RGBA2YUV_YV12) { performTest(4, 1, CVTCODE(RGBA2YUV_YV12), 1); } +OCL_TEST_P(CvtColor_RGB2YUV_420, BGRA2YUV_YV12) { performTest(4, 1, CVTCODE(BGRA2YUV_YV12), 1); } +OCL_TEST_P(CvtColor_RGB2YUV_420, RGB2YUV_YV12) { performTest(3, 1, CVTCODE(RGB2YUV_YV12), 1); } +OCL_TEST_P(CvtColor_RGB2YUV_420, BGR2YUV_YV12) { performTest(3, 1, CVTCODE(BGR2YUV_YV12), 1); } +OCL_TEST_P(CvtColor_RGB2YUV_420, RGBA2YUV_IYUV) { performTest(4, 1, CVTCODE(RGBA2YUV_IYUV), 1); } +OCL_TEST_P(CvtColor_RGB2YUV_420, BGRA2YUV_IYUV) { performTest(4, 1, CVTCODE(BGRA2YUV_IYUV), 1); } +OCL_TEST_P(CvtColor_RGB2YUV_420, RGB2YUV_IYUV) { performTest(3, 1, CVTCODE(RGB2YUV_IYUV), 1); } +OCL_TEST_P(CvtColor_RGB2YUV_420, BGR2YUV_IYUV) { performTest(3, 1, CVTCODE(BGR2YUV_IYUV), 1); } // YUV422 -> RGBA From a8aa6381d96802bae18110d9c324597059d4db1d Mon Sep 17 00:00:00 2001 From: vbystricky Date: Mon, 20 Oct 2014 17:43:18 +0400 Subject: [PATCH 10/26] Optimize OpenCL version of conversScaleAbs function --- modules/core/include/opencv2/core/ocl.hpp | 2 +- modules/core/src/convert.cpp | 20 ++++++++++++++++---- modules/core/src/ocl.cpp | 4 +++- 3 files changed, 20 insertions(+), 6 deletions(-) diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp index a84b98d151..0ee492e097 100644 --- a/modules/core/include/opencv2/core/ocl.hpp +++ b/modules/core/include/opencv2/core/ocl.hpp @@ -618,7 +618,7 @@ CV_EXPORTS int predictOptimalVectorWidth(InputArray src1, InputArray src2 = noAr InputArray src7 = noArray(), InputArray src8 = noArray(), InputArray src9 = noArray(), OclVectorStrategy strat = OCL_VECTOR_DEFAULT); -CV_EXPORTS int checkOptimalVectorWidth(int *vectorWidths, +CV_EXPORTS int checkOptimalVectorWidth(const int *vectorWidths, InputArray src1, InputArray src2 = noArray(), InputArray src3 = noArray(), InputArray src4 = noArray(), InputArray src5 = noArray(), InputArray src6 = noArray(), InputArray src7 = noArray(), InputArray src8 = noArray(), InputArray src9 = noArray(), diff --git a/modules/core/src/convert.cpp b/modules/core/src/convert.cpp index d6da1a2253..829b984c9f 100644 --- a/modules/core/src/convert.cpp +++ b/modules/core/src/convert.cpp @@ -3275,13 +3275,26 @@ static BinaryFunc getConvertScaleFunc(int sdepth, int ddepth) static bool ocl_convertScaleAbs( InputArray _src, OutputArray _dst, double alpha, double beta ) { const ocl::Device & d = ocl::Device::getDefault(); - int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), - kercn = ocl::predictOptimalVectorWidth(_src, _dst), rowsPerWI = d.isIntel() ? 4 : 1; - bool doubleSupport = d.doubleFPConfig() > 0; + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + bool doubleSupport = d.doubleFPConfig() > 0; if (!doubleSupport && depth == CV_64F) return false; + _dst.create(_src.size(), CV_8UC(cn)); + int kercn = 1; + if (d.isIntel()) + { + static const int vectorWidths[] = {4, 4, 4, 4, 4, 4, 4, -1}; + kercn = ocl::checkOptimalVectorWidth( vectorWidths, _src, _dst, + noArray(), noArray(), noArray(), + noArray(), noArray(), noArray(), + noArray(), ocl::OCL_VECTOR_MAX); + } + else + kercn = ocl::predictOptimalVectorWidthMax(_src, _dst); + + int rowsPerWI = d.isIntel() ? 4 : 1; char cvt[2][50]; int wdepth = std::max(depth, CV_32F); String build_opt = format("-D OP_CONVERT_SCALE_ABS -D UNARY_OP -D dstT=%s -D srcT1=%s" @@ -3299,7 +3312,6 @@ static bool ocl_convertScaleAbs( InputArray _src, OutputArray _dst, double alpha return false; UMat src = _src.getUMat(); - _dst.create(src.size(), CV_8UC(cn)); UMat dst = _dst.getUMat(); ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src), diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index c333b08c40..faf45d3c87 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -4531,12 +4531,14 @@ int predictOptimalVectorWidth(InputArray src1, InputArray src2, InputArray src3, return checkOptimalVectorWidth(vectorWidths, src1, src2, src3, src4, src5, src6, src7, src8, src9, strat); } -int checkOptimalVectorWidth(int *vectorWidths, +int checkOptimalVectorWidth(const int *vectorWidths, InputArray src1, InputArray src2, InputArray src3, InputArray src4, InputArray src5, InputArray src6, InputArray src7, InputArray src8, InputArray src9, OclVectorStrategy strat) { + CV_Assert(vectorWidths); + int ref_type = src1.type(); std::vector offsets, steps, cols; From 62d4393883b21b55b465ace2df232405c66c22f2 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Wed, 22 Oct 2014 17:43:29 +0300 Subject: [PATCH 11/26] Small optimization for buildMaps kernels in case rowsPerWI > 1 --- modules/stitching/src/opencl/warpers.cl | 39 +++++++++++-------------- modules/stitching/src/warpers.cpp | 6 ++-- 2 files changed, 20 insertions(+), 25 deletions(-) diff --git a/modules/stitching/src/opencl/warpers.cl b/modules/stitching/src/opencl/warpers.cl index 9b5619fcad..7ec87ae2f8 100644 --- a/modules/stitching/src/opencl/warpers.cl +++ b/modules/stitching/src/opencl/warpers.cl @@ -56,19 +56,19 @@ __kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xm int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset)); int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset)); + float u = tl_u + du; + float x_ = u * scale - ct[0]; + float ct1 = 1 - ct[2]; + for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step, ymap_index += ymap_step) { __global float * xmap = (__global float *)(xmapptr + xmap_index); __global float * ymap = (__global float *)(ymapptr + ymap_index); - float u = tl_u + du; float v = tl_v + dv; + float y_ = v * scale - ct[1]; - float x_ = u / scale - ct[0]; - float y_ = v / scale - ct[1]; - - float ct1 = 1 - ct[2]; float x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * ct1)); float y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * ct1)); float z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * ct1)); @@ -94,22 +94,19 @@ __kernel void buildWarpCylindricalMaps(__global uchar * xmapptr, int xmap_step, int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset)); int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset)); + float u = (tl_u + du) * scale; + float x_, z_; + x_ = sincos(u, &z_); + for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step, ymap_index += ymap_step) { __global float * xmap = (__global float *)(xmapptr + xmap_index); __global float * ymap = (__global float *)(ymapptr + ymap_index); - float u = tl_u + du; - float v = tl_v + dv; - float x, y; + float y_ = (tl_v + dv) * scale; - u /= scale; - float x_, y_, z_; - x_ = sincos(u, &z_); - y_ = v / scale; - - float z; + float x, y, z; x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * z_)); y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * z_)); z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * z_)); @@ -137,25 +134,23 @@ __kernel void buildWarpSphericalMaps(__global uchar * xmapptr, int xmap_step, in int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset)); int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset)); + float u = (tl_u + du) * scale; + float cosu, sinu = sincos(u, &cosu); + for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step, ymap_index += ymap_step) { __global float * xmap = (__global float *)(xmapptr + xmap_index); __global float * ymap = (__global float *)(ymapptr + ymap_index); - float u = tl_u + du; - float v = tl_v + dv; - float x, y; + float v = (tl_v + dv) * scale; - v /= scale; - u /= scale; - - float cosv, sinv = sincos(v, &cosv), cosu, sinu = sincos(u, &cosu); + float cosv, sinv = sincos(v, &cosv); float x_ = sinv * sinu; float y_ = -cosv; float z_ = sinv * cosu; - float z; + float x, y, z; x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * z_)); y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * z_)); z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * z_)); diff --git a/modules/stitching/src/warpers.cpp b/modules/stitching/src/warpers.cpp index f474010ac5..2711a012f6 100644 --- a/modules/stitching/src/warpers.cpp +++ b/modules/stitching/src/warpers.cpp @@ -110,7 +110,7 @@ Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, InputArra k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), ocl::KernelArg::PtrReadOnly(uk_rinv), ocl::KernelArg::PtrReadOnly(ut), - dst_tl.x, dst_tl.y, projector_.scale, rowsPerWI); + dst_tl.x, dst_tl.y, 1/projector_.scale, rowsPerWI); size_t globalsize[2] = { dsize.width, (dsize.height + rowsPerWI - 1) / rowsPerWI }; if (k.run(2, globalsize, NULL, true)) @@ -388,7 +388,7 @@ Rect SphericalWarper::buildMaps(Size src_size, InputArray K, InputArray R, Outpu UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), uk_rinv = k_rinv.getUMat(ACCESS_READ); k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), - ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, projector_.scale, rowsPerWI); + ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, 1/projector_.scale, rowsPerWI); size_t globalsize[2] = { dsize.width, (dsize.height + rowsPerWI - 1) / rowsPerWI }; if (k.run(2, globalsize, NULL, true)) @@ -436,7 +436,7 @@ Rect CylindricalWarper::buildMaps(Size src_size, InputArray K, InputArray R, Out UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), uk_rinv = k_rinv.getUMat(ACCESS_READ); k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), - ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, projector_.scale, + ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, 1/projector_.scale, rowsPerWI); size_t globalsize[2] = { dsize.width, (dsize.height + rowsPerWI - 1) / rowsPerWI }; From 3a263c6326d2e46ef85474a0eb991160d483031c Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Thu, 23 Oct 2014 14:23:37 +0300 Subject: [PATCH 12/26] Added tests for Image2D --- modules/core/src/ocl.cpp | 3 ++ modules/core/test/ocl/test_image2d.cpp | 43 ++++++++++++++++++++++++++ 2 files changed, 46 insertions(+) create mode 100644 modules/core/test/ocl/test_image2d.cpp diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index c333b08c40..708730ee50 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -4624,6 +4624,9 @@ struct Image2D::Impl static bool isFormatSupported(cl_image_format format) { + if (!haveOpenCL()) + CV_Error(Error::OpenCLApiCallError, "OpenCL runtime not found!"); + cl_context context = (cl_context)Context::getDefault().ptr(); // Figure out how many formats are supported by this context. cl_uint numFormats = 0; diff --git a/modules/core/test/ocl/test_image2d.cpp b/modules/core/test/ocl/test_image2d.cpp new file mode 100644 index 0000000000..412761a7f8 --- /dev/null +++ b/modules/core/test/ocl/test_image2d.cpp @@ -0,0 +1,43 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. + +// Copyright (C) 2014, Itseez, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. + +#include "../test_precomp.hpp" +#include "opencv2/ts/ocl_test.hpp" + +#ifdef HAVE_OPENCL + +namespace cvtest { +namespace ocl { + +PARAM_TEST_CASE(Image2DBasicTest, int, int) +{ + int depth, ch; + + +}; + +TEST(Image2D, turnOffOpenCL) +{ + if (cv::ocl::haveOpenCL()) + { + // save the current state + bool useOCL = cv::ocl::useOpenCL(); + + cv::ocl::setUseOpenCL(true); + UMat um(128, 128, CV_8UC1); + + cv::ocl::setUseOpenCL(false); + cv::ocl::Image2D image(um); + + // reset state to the previous one + cv::ocl::setUseOpenCL(useOCL); + } +} + +} } // namespace cvtest::ocl + +#endif // HAVE_OPENCL \ No newline at end of file From 21ad8e92d4b039817e4004bfc952be0576eb4d1d Mon Sep 17 00:00:00 2001 From: vbystricky Date: Mon, 29 Sep 2014 12:46:56 +0400 Subject: [PATCH 13/26] Optimize OpenCL version of StereoBM function Fix problems on NVidia devices. --- modules/calib3d/src/opencl/stereobm.cl | 267 ++++++++++++++----------- modules/calib3d/src/stereobm.cpp | 33 ++- 2 files changed, 167 insertions(+), 133 deletions(-) diff --git a/modules/calib3d/src/opencl/stereobm.cl b/modules/calib3d/src/opencl/stereobm.cl index 73402a6a1d..86e4c24d33 100644 --- a/modules/calib3d/src/opencl/stereobm.cl +++ b/modules/calib3d/src/opencl/stereobm.cl @@ -44,217 +44,251 @@ ////////////////////////////////////////// stereoBM ////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////////////////////// -#ifdef csize - #define MAX_VAL 32767 -void calcDisp(__local short * cost, __global short * disp, int uniquenessRatio, int mindisp, int ndisp, int w, - __local int * bestDisp, __local int * bestCost, int d, int x, int y, int cols, int rows, int wsz2) +#ifndef WSZ +#define WSZ 2 +#endif + +#define WSZ2 (WSZ / 2) + +#ifdef DEFINE_KERNEL_STEREOBM + +#define DISPARITY_SHIFT 4 +#define FILTERED ((MIN_DISP - 1) << DISPARITY_SHIFT) + +void calcDisp(__local short * cost, __global short * disp, int uniquenessRatio, + __local int * bestDisp, __local int * bestCost, int d, int x, int y, int cols, int rows) { - short FILTERED = (mindisp - 1)<<4; - int best_disp = *bestDisp, best_cost = *bestCost, best_disp_back = ndisp - best_disp - 1; + int best_disp = *bestDisp, best_cost = *bestCost; + barrier(CLK_LOCAL_MEM_FENCE); short c = cost[0]; + int thresh = best_cost + (best_cost * uniquenessRatio / 100); + bool notUniq = ( (c <= thresh) && (d < (best_disp - 1) || d > (best_disp + 1) ) ); - int thresh = best_cost + (best_cost * uniquenessRatio/100); - bool notUniq = ( (c <= thresh) && (d < (best_disp_back - 1) || d > (best_disp_back + 1) ) ); - - if(notUniq) + if (notUniq) *bestCost = FILTERED; barrier(CLK_LOCAL_MEM_FENCE); - if( *bestCost != FILTERED && x < cols-wsz2-mindisp && y < rows-wsz2 && d == best_disp_back) + if( *bestCost != FILTERED && x < cols - WSZ2 - MIN_DISP && y < rows - WSZ2 && d == best_disp) { - int y3 = (best_disp_back > 0) ? cost[-w] : cost[w], - y2 = c, - y1 = (best_disp_back < ndisp-1) ? cost[w] : cost[-w]; - int d_aprox = y3+y1-2*y2 + abs(y3-y1); - disp[0] = (short)(((best_disp_back + mindisp)*256 + (d_aprox != 0 ? (y3-y1)*256/d_aprox : 0) + 15) >> 4); + int d_aprox = 0; + int yp =0, yn = 0; + if ((0 < best_disp) && (best_disp < NUM_DISP - 1)) + { + yp = cost[-2 * BLOCK_SIZE_Y]; + yn = cost[2 * BLOCK_SIZE_Y]; + d_aprox = yp + yn - 2 * c + abs(yp - yn); + } + disp[0] = (short)(((best_disp + MIN_DISP)*256 + (d_aprox != 0 ? (yp - yn) * 256 / d_aprox : 0) + 15) >> 4); } } -int calcLocalIdx(int x, int y, int d, int w) -{ - return d*2*w + (w - 1 - y + x); -} - -void calcNewCoordinates(int * x, int * y, int nthread) -{ - int oldX = *x - (1-nthread), oldY = *y; - *x = (oldX == oldY) ? (0*nthread + (oldX + 2)*(1-nthread) ) : (oldX+1)*(1-nthread) + (oldX+1)*nthread; - *y = (oldX == oldY) ? (0*(1-nthread) + (oldY + 1)*nthread) : oldY + 1*(1-nthread); -} - short calcCostBorder(__global const uchar * leftptr, __global const uchar * rightptr, int x, int y, int nthread, - int wsz2, short * costbuf, int * h, int cols, int d, short cost, int winsize) + short * costbuf, int *h, int cols, int d, short cost) { - int head = (*h)%wsz; + int head = (*h) % WSZ; __global const uchar * left, * right; - int idx = mad24(y+wsz2*(2*nthread-1), cols, x+wsz2*(1-2*nthread)); + int idx = mad24(y + WSZ2 * (2 * nthread - 1), cols, x + WSZ2 * (1 - 2 * nthread)); left = leftptr + idx; right = rightptr + (idx - d); - int shift = 1*nthread + cols*(1-nthread); short costdiff = 0; - for(int i = 0; i < winsize; i++) + if (0 == nthread) { - costdiff += abs( left[0] - right[0] ); - left += shift; - right += shift; + #pragma unroll + for (int i = 0; i < WSZ; i++) + { + costdiff += abs( left[0] - right[0] ); + left += cols; + right += cols; + } + } + else // (1 == nthread) + { + #pragma unroll + for (int i = 0; i < WSZ; i++) + { + costdiff += abs(left[i] - right[i]); + } } cost += costdiff - costbuf[head]; costbuf[head] = costdiff; - (*h) = (*h)%wsz + 1; + *h = head + 1; return cost; } short calcCostInside(__global const uchar * leftptr, __global const uchar * rightptr, int x, int y, - int wsz2, int cols, int d, short cost_up_left, short cost_up, short cost_left, - int winsize) + int cols, int d, short cost_up_left, short cost_up, short cost_left) { __global const uchar * left, * right; - int idx = mad24(y-wsz2-1, cols, x-wsz2-1); + int idx = mad24(y - WSZ2 - 1, cols, x - WSZ2 - 1); left = leftptr + idx; right = rightptr + (idx - d); - int idx2 = winsize*cols; + int idx2 = WSZ*cols; uchar corrner1 = abs(left[0] - right[0]), - corrner2 = abs(left[winsize] - right[winsize]), + corrner2 = abs(left[WSZ] - right[WSZ]), corrner3 = abs(left[idx2] - right[idx2]), - corrner4 = abs(left[idx2 + winsize] - right[idx2 + winsize]); + corrner4 = abs(left[idx2 + WSZ] - right[idx2 + WSZ]); return cost_up + cost_left - cost_up_left + corrner1 - corrner2 - corrner3 + corrner4; } -__kernel void stereoBM(__global const uchar * leftptr, __global const uchar * rightptr, __global uchar * dispptr, - int disp_step, int disp_offset, int rows, int cols, int mindisp, int ndisp, - int preFilterCap, int textureTreshold, int uniquenessRatio, int sizeX, int sizeY, int winsize) +__kernel void stereoBM(__global const uchar * leftptr, + __global const uchar * rightptr, + __global uchar * dispptr, int disp_step, int disp_offset, + int rows, int cols, // rows, cols of left and right images, not disp + int textureTreshold, int uniquenessRatio) { - int gx = get_global_id(0)*sizeX; - int gy = get_global_id(1)*sizeY; - int lz = get_local_id(2); + int lz = get_local_id(0); + int gx = get_global_id(1) * BLOCK_SIZE_X; + int gy = get_global_id(2) * BLOCK_SIZE_Y; - int nthread = lz/ndisp; - int d = lz%ndisp; - int wsz2 = wsz/2; + int nthread = lz / NUM_DISP; + int disp_idx = lz % NUM_DISP; __global short * disp; __global const uchar * left, * right; - __local short costFunc[csize]; + __local short costFunc[2 * BLOCK_SIZE_Y * NUM_DISP]; + __local short * cost; __local int best_disp[2]; __local int best_cost[2]; best_cost[nthread] = MAX_VAL; - best_disp[nthread] = MAX_VAL; + best_disp[nthread] = -1; barrier(CLK_LOCAL_MEM_FENCE); - short costbuf[wsz]; + short costbuf[WSZ]; int head = 0; - int shiftX = wsz2 + ndisp + mindisp - 1; - int shiftY = wsz2; + int shiftX = WSZ2 + NUM_DISP + MIN_DISP - 1; + int shiftY = WSZ2; int x = gx + shiftX, y = gy + shiftY, lx = 0, ly = 0; - int costIdx = calcLocalIdx(lx, ly, d, sizeY); + int costIdx = disp_idx * 2 * BLOCK_SIZE_Y + (BLOCK_SIZE_Y - 1); cost = costFunc + costIdx; int tempcost = 0; - if(x < cols-wsz2-mindisp && y < rows-wsz2) + if (x < cols - WSZ2 - MIN_DISP && y < rows - WSZ2) { - int shift = 1*nthread + cols*(1-nthread); - for(int i = 0; i < winsize; i++) + if (0 == nthread) { - int idx = mad24(y-wsz2+i*nthread, cols, x-wsz2+i*(1-nthread)); - left = leftptr + idx; - right = rightptr + (idx - d); - short costdiff = 0; - for(int j = 0; j < winsize; j++) + #pragma unroll + for (int i = 0; i < WSZ; i++) { - costdiff += abs( left[0] - right[0] ); - left += shift; - right += shift; + int idx = mad24(y - WSZ2, cols, x - WSZ2 + i); + left = leftptr + idx; + right = rightptr + (idx - disp_idx); + short costdiff = 0; + for(int j = 0; j < WSZ; j++) + { + costdiff += abs( left[0] - right[0] ); + left += cols; + right += cols; + } + costbuf[i] = costdiff; } - if(nthread==1) + } + else // (1 == nthread) + { + #pragma unroll + for (int i = 0; i < WSZ; i++) { + int idx = mad24(y - WSZ2 + i, cols, x - WSZ2); + left = leftptr + idx; + right = rightptr + (idx - disp_idx); + short costdiff = 0; + for (int j = 0; j < WSZ; j++) + { + costdiff += abs( left[j] - right[j]); + } tempcost += costdiff; + costbuf[i] = costdiff; } - costbuf[head] = costdiff; - head++; } } - if(nthread==1) + if (nthread == 1) { cost[0] = tempcost; - atomic_min(best_cost+nthread, tempcost); + atomic_min(best_cost + 1, tempcost); } barrier(CLK_LOCAL_MEM_FENCE); - if(best_cost[1] == tempcost) - atomic_min(best_disp + 1, ndisp - d - 1); + if (best_cost[1] == tempcost) + atomic_max(best_disp + 1, disp_idx); barrier(CLK_LOCAL_MEM_FENCE); - int dispIdx = mad24(gy, disp_step, disp_offset + gx*(int)sizeof(short)); + int dispIdx = mad24(gy, disp_step, mad24((int)sizeof(short), gx, disp_offset)); disp = (__global short *)(dispptr + dispIdx); - calcDisp(cost, disp, uniquenessRatio, mindisp, ndisp, 2*sizeY, - best_disp + 1, best_cost+1, d, x, y, cols, rows, wsz2); + calcDisp(cost, disp, uniquenessRatio, best_disp + 1, best_cost + 1, disp_idx, x, y, cols, rows); barrier(CLK_LOCAL_MEM_FENCE); lx = 1 - nthread; ly = nthread; - for(int i = 0; i < sizeY*sizeX/2; i++) + for (int i = 0; i < BLOCK_SIZE_Y * BLOCK_SIZE_X / 2; i++) { - x = (lx < sizeX) ? gx + shiftX + lx : cols; - y = (ly < sizeY) ? gy + shiftY + ly : rows; + x = (lx < BLOCK_SIZE_X) ? gx + shiftX + lx : cols; + y = (ly < BLOCK_SIZE_Y) ? gy + shiftY + ly : rows; best_cost[nthread] = MAX_VAL; - best_disp[nthread] = MAX_VAL; + best_disp[nthread] = -1; barrier(CLK_LOCAL_MEM_FENCE); - costIdx = calcLocalIdx(lx, ly, d, sizeY); + costIdx = mad24(2 * BLOCK_SIZE_Y, disp_idx, (BLOCK_SIZE_Y - 1 - ly + lx)); + if (0 > costIdx) + costIdx = BLOCK_SIZE_Y - 1; cost = costFunc + costIdx; - - if(x < cols-wsz2-mindisp && y < rows-wsz2 ) + if (x < cols - WSZ2 - MIN_DISP && y < rows - WSZ2) { - tempcost = ( ly*(1-nthread) + lx*nthread == 0 ) ? - calcCostBorder(leftptr, rightptr, x, y, nthread, wsz2, costbuf, &head, cols, d, - cost[2*nthread-1], winsize) : - calcCostInside(leftptr, rightptr, x, y, wsz2, cols, d, - cost[0], cost[1], cost[-1], winsize); + tempcost = (ly * (1 - nthread) + lx * nthread == 0) ? + calcCostBorder(leftptr, rightptr, x, y, nthread, costbuf, &head, cols, disp_idx, cost[2*nthread-1]) : + calcCostInside(leftptr, rightptr, x, y, cols, disp_idx, cost[0], cost[1], cost[-1]); } cost[0] = tempcost; atomic_min(best_cost + nthread, tempcost); barrier(CLK_LOCAL_MEM_FENCE); - if(best_cost[nthread] == tempcost) - atomic_min(best_disp + nthread, ndisp - d - 1); + if (best_cost[nthread] == tempcost) + atomic_max(best_disp + nthread, disp_idx); barrier(CLK_LOCAL_MEM_FENCE); - int dispIdx = mad24(gy+ly, disp_step, disp_offset + (gx+lx)*(int)sizeof(short)); + dispIdx = mad24(gy + ly, disp_step, mad24((int)sizeof(short), (gx + lx), disp_offset)); disp = (__global short *)(dispptr + dispIdx); - calcDisp(cost, disp, uniquenessRatio, mindisp, ndisp, 2*sizeY, - best_disp + nthread, best_cost + nthread, d, x, y, cols, rows, wsz2); + calcDisp(cost, disp, uniquenessRatio, best_disp + nthread, best_cost + nthread, disp_idx, x, y, cols, rows); + barrier(CLK_LOCAL_MEM_FENCE); - calcNewCoordinates(&lx, &ly, nthread); + if (lx + nthread - 1 == ly) + { + lx = (lx + nthread + 1) * (1 - nthread); + ly = (ly + 1) * nthread; + } + else + { + lx += nthread; + ly = ly - nthread + 1; + } } } - -#endif +#endif //DEFINE_KERNEL_STEREOBM ////////////////////////////////////////////////////////////////////////////////////////////////// /////////////////////////////////////// Norm Prefiler //////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////////////////////// __kernel void prefilter_norm(__global unsigned char *input, __global unsigned char *output, - int rows, int cols, int prefilterCap, int winsize, int scale_g, int scale_s) + int rows, int cols, int prefilterCap, int scale_g, int scale_s) { + // prefilterCap in range 1..63, checked in StereoBMImpl::compute + int x = get_global_id(0); int y = get_global_id(1); - int wsz2 = winsize/2; if(x < cols && y < rows) { @@ -262,13 +296,13 @@ __kernel void prefilter_norm(__global unsigned char *input, __global unsigned ch input[y * cols + max(x-1,0)] * 1 + input[ y * cols + x] * 4 + input[y * cols + min(x+1, cols-1)] * 1 + input[min(y+1, rows-1) * cols + x] * 1; int cov2 = 0; - for(int i = -wsz2; i < wsz2+1; i++) - for(int j = -wsz2; j < wsz2+1; j++) + for(int i = -WSZ2; i < WSZ2+1; i++) + for(int j = -WSZ2; j < WSZ2+1; j++) cov2 += input[clamp(y+i, 0, rows-1) * cols + clamp(x+j, 0, cols-1)]; int res = (cov1*scale_g - cov2*scale_s)>>10; - res = min(clamp(res, -prefilterCap, prefilterCap) + prefilterCap, 255); - output[y * cols + x] = res & 0xFF; + res = clamp(res, -prefilterCap, prefilterCap) + prefilterCap; + output[y * cols + x] = res; } } @@ -280,20 +314,21 @@ __kernel void prefilter_norm(__global unsigned char *input, __global unsigned ch __kernel void prefilter_xsobel(__global unsigned char *input, __global unsigned char *output, int rows, int cols, int prefilterCap) { + // prefilterCap in range 1..63, checked in StereoBMImpl::compute int x = get_global_id(0); int y = get_global_id(1); if(x < cols && y < rows) { - output[y * cols + x] = min(prefilterCap, 255) & 0xFF; - } + if (0 < x && !((y == rows-1) & (rows%2==1) ) ) + { + int cov = input[ ((y > 0) ? y-1 : y+1) * cols + (x-1)] * (-1) + input[ ((y > 0) ? y-1 : y+1) * cols + ((x 0 && !((y == rows-1)&(rows%2==1) ) ) - { - int cov = input[ ((y > 0) ? y-1 : y+1) * cols + (x-1)] * (-1) + input[ ((y > 0) ? y-1 : y+1) * cols + ((xSADWindowSize; int wsz2 = wsz/2; - int sizeX = std::max(11, 27 - ocl::Device::getDefault().maxComputeUnits() ), sizeY = sizeX-1, N = ndisp*2; + ocl::Device devDef = ocl::Device::getDefault(); + int sizeX = devDef.isIntel() ? 32 : std::max(11, 27 - devDef.maxComputeUnits()), + sizeY = sizeX - 1, + N = ndisp * 2; - ocl::Kernel k("stereoBM", ocl::calib3d::stereobm_oclsrc, cv::format("-D csize=%d -D wsz=%d", (2*sizeY)*ndisp, wsz) ); + cv::String opt = cv::format("-D DEFINE_KERNEL_STEREOBM -D MIN_DISP=%d -D NUM_DISP=%d" + " -D BLOCK_SIZE_X=%d -D BLOCK_SIZE_Y=%d -D WSZ=%d", + mindisp, ndisp, + sizeX, sizeY, wsz); + ocl::Kernel k("stereoBM", ocl::calib3d::stereobm_oclsrc, opt); if(k.empty()) return false; @@ -753,15 +760,14 @@ static bool ocl_stereobm( InputArray _left, InputArray _right, int cols = left.cols, rows = left.rows; _disp.create(_left.size(), CV_16S); - _disp.setTo((mindisp - 1)<<4); + _disp.setTo((mindisp - 1) << 4); Rect roi = Rect(Point(wsz2 + mindisp + ndisp - 1, wsz2), Point(cols-wsz2-mindisp, rows-wsz2) ); UMat disp = (_disp.getUMat())(roi); - int globalX = disp.cols/sizeX, globalY = disp.rows/sizeY; - globalX += (disp.cols%sizeX) > 0 ? 1 : 0; - globalY += (disp.rows%sizeY) > 0 ? 1 : 0; - size_t globalThreads[3] = { globalX, globalY, N}; - size_t localThreads[3] = {1, 1, N}; + int globalX = (disp.cols + sizeX - 1) / sizeX, + globalY = (disp.rows + sizeY - 1) / sizeY; + size_t globalThreads[3] = {N, globalX, globalY}; + size_t localThreads[3] = {N, 1, 1}; int idx = 0; idx = k.set(idx, ocl::KernelArg::PtrReadOnly(left)); @@ -769,15 +775,8 @@ static bool ocl_stereobm( InputArray _left, InputArray _right, idx = k.set(idx, ocl::KernelArg::WriteOnlyNoSize(disp)); idx = k.set(idx, rows); idx = k.set(idx, cols); - idx = k.set(idx, mindisp); - idx = k.set(idx, ndisp); - idx = k.set(idx, state->preFilterCap); idx = k.set(idx, state->textureThreshold); idx = k.set(idx, state->uniquenessRatio); - idx = k.set(idx, sizeX); - idx = k.set(idx, sizeY); - idx = k.set(idx, wsz); - return k.run(3, globalThreads, localThreads, false); } From 579a7fff6d557d6af4ab5816dc3a098707aa3c33 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Fri, 24 Oct 2014 14:29:38 +0400 Subject: [PATCH 14/26] ocl: restore clFinish() in unmap() for AMD devices This reverts commit 7d91b8efcd053304e5b0fb5afff06e151b2d4096. --- modules/core/src/ocl.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index c333b08c40..79a5367b33 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -3979,6 +3979,11 @@ public: u->markDeviceMemMapped(false); CV_Assert( (retval = clEnqueueUnmapMemObject(q, (cl_mem)u->handle, u->data, 0, 0, 0)) == CL_SUCCESS ); + if (Device::getDefault().isAMD()) + { + // required for multithreaded applications (see stitching test) + CV_OclDbgAssert(clFinish(q) == CL_SUCCESS); + } u->data = 0; } else if( u->copyOnMap() && u->deviceCopyObsolete() ) From 237cb9314378f0819e10a9c95eb41ec61e2c3ebe Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Fri, 24 Oct 2014 13:55:16 +0300 Subject: [PATCH 15/26] Added extra checks to ocl::Image2D --- modules/core/src/ocl.cpp | 9 +++- modules/core/test/ocl/test_image2d.cpp | 61 ++++++++++++++++++++++++-- 2 files changed, 65 insertions(+), 5 deletions(-) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 708730ee50..9200c29f55 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -4650,6 +4650,10 @@ struct Image2D::Impl void init(const UMat &src, bool norm, bool alias) { + if (!haveOpenCL()) + CV_Error(Error::OpenCLApiCallError, "OpenCL runtime not found!"); + + CV_Assert(!src.empty()); CV_Assert(ocl::Device::getDefault().imageSupport()); int err, depth = src.depth(), cn = src.channels(); @@ -4659,6 +4663,9 @@ struct Image2D::Impl if (!isFormatSupported(format)) CV_Error(Error::OpenCLApiCallError, "Image format is not supported"); + if (alias && !src.handle(ACCESS_RW)) + CV_Error(Error::OpenCLApiCallError, "Incorrect UMat, handle is null"); + cl_context context = (cl_context)Context::getDefault().ptr(); cl_command_queue queue = (cl_command_queue)Queue::getDefault().ptr(); @@ -4743,7 +4750,7 @@ bool Image2D::canCreateAlias(const UMat &m) { bool ret = false; const Device & d = ocl::Device::getDefault(); - if (d.imageFromBufferSupport()) + if (d.imageFromBufferSupport() && !m.empty()) { // This is the required pitch alignment in pixels uint pitchAlign = d.imagePitchAlignment(); diff --git a/modules/core/test/ocl/test_image2d.cpp b/modules/core/test/ocl/test_image2d.cpp index 412761a7f8..dcfc701f11 100644 --- a/modules/core/test/ocl/test_image2d.cpp +++ b/modules/core/test/ocl/test_image2d.cpp @@ -13,12 +13,55 @@ namespace cvtest { namespace ocl { -PARAM_TEST_CASE(Image2DBasicTest, int, int) +TEST(Image2D, createAliasEmptyUMat) { - int depth, ch; + if (cv::ocl::haveOpenCL()) + { + UMat um; + EXPECT_FALSE(cv::ocl::Image2D::canCreateAlias(um)); + } + else + std::cout << "OpenCL runtime not found. Test skipped." << std::endl; +} +TEST(Image2D, createImage2DWithEmptyUMat) +{ + if (cv::ocl::haveOpenCL()) + { + UMat um; + EXPECT_ANY_THROW(cv::ocl::Image2D image(um)); + } + else + std::cout << "OpenCL runtime not found. Test skipped." << std::endl; +} -}; +TEST(Image2D, createAlias) +{ + if (cv::ocl::haveOpenCL()) + { + const cv::ocl::Device & d = cv::ocl::Device::getDefault(); + int minor = d.deviceVersionMinor(), major = d.deviceVersionMajor(); + + // aliases is OpenCL 1.2 extension + if (1 < major || (1 == major && 2 <= minor)) + { + UMat um(128, 128, CV_8UC1); + bool isFormatSupported = false, canCreateAlias = false; + + EXPECT_NO_THROW(isFormatSupported = cv::ocl::Image2D::isFormatSupported(CV_8U, 1, false)); + EXPECT_NO_THROW(canCreateAlias = cv::ocl::Image2D::canCreateAlias(um)); + + if (isFormatSupported && canCreateAlias) + { + EXPECT_NO_THROW(cv::ocl::Image2D image(um, false, true)); + } + else + std::cout << "Impossible to create alias for selected image. Test skipped." << std::endl; + } + } + else + std::cout << "OpenCL runtime not found. Test skipped" << std::endl; +} TEST(Image2D, turnOffOpenCL) { @@ -26,16 +69,26 @@ TEST(Image2D, turnOffOpenCL) { // save the current state bool useOCL = cv::ocl::useOpenCL(); + bool isFormatSupported = false; cv::ocl::setUseOpenCL(true); UMat um(128, 128, CV_8UC1); cv::ocl::setUseOpenCL(false); - cv::ocl::Image2D image(um); + EXPECT_NO_THROW(isFormatSupported = cv::ocl::Image2D::isFormatSupported(CV_8U, 1, true)); + + if (isFormatSupported) + { + EXPECT_NO_THROW(cv::ocl::Image2D image(um)); + } + else + std::cout << "CV_8UC1 is not supported for OpenCL images. Test skipped." << std::endl; // reset state to the previous one cv::ocl::setUseOpenCL(useOCL); } + else + std::cout << "OpenCL runtime not found. Test skipped." << std::endl; } } } // namespace cvtest::ocl From 8466911ad04f2835c41edcaa1bb3e5da117ad611 Mon Sep 17 00:00:00 2001 From: vbystricky Date: Fri, 24 Oct 2014 14:05:39 +0400 Subject: [PATCH 16/26] Move _dst.create() to the begining of scaleAdd function --- modules/core/src/matmul.cpp | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/modules/core/src/matmul.cpp b/modules/core/src/matmul.cpp index d171dc3da7..e3e1720ef0 100644 --- a/modules/core/src/matmul.cpp +++ b/modules/core/src/matmul.cpp @@ -2173,14 +2173,18 @@ typedef void (*ScaleAddFunc)(const uchar* src1, const uchar* src2, uchar* dst, i static bool ocl_scaleAdd( InputArray _src1, double alpha, InputArray _src2, OutputArray _dst, int type ) { const ocl::Device & d = ocl::Device::getDefault(); - int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), wdepth = std::max(depth, CV_32F), - kercn = ocl::predictOptimalVectorWidth(_src1, _src2, _dst), rowsPerWI = d.isIntel() ? 4 : 1; + bool doubleSupport = d.doubleFPConfig() > 0; Size size = _src1.size(); - + int depth = CV_MAT_DEPTH(type); if ( (!doubleSupport && depth == CV_64F) || size != _src2.size() ) return false; + _dst.create(size, type); + int cn = CV_MAT_CN(type), wdepth = std::max(depth, CV_32F); + int kercn = ocl::predictOptimalVectorWidthMax(_src1, _src2, _dst), + rowsPerWI = d.isIntel() ? 4 : 1; + char cvt[2][50]; ocl::Kernel k("KF", ocl::core::arithm_oclsrc, format("-D OP_SCALE_ADD -D BINARY_OP -D dstT=%s -D workT=%s -D convertToWT1=%s" @@ -2195,9 +2199,7 @@ static bool ocl_scaleAdd( InputArray _src1, double alpha, InputArray _src2, Outp if (k.empty()) return false; - UMat src1 = _src1.getUMat(), src2 = _src2.getUMat(); - _dst.create(size, type); - UMat dst = _dst.getUMat(); + UMat src1 = _src1.getUMat(), src2 = _src2.getUMat(), dst = _dst.getUMat(); ocl::KernelArg src1arg = ocl::KernelArg::ReadOnlyNoSize(src1), src2arg = ocl::KernelArg::ReadOnlyNoSize(src2), From d78bc3c3219bd52a9ae99177e300267ff49a3269 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Fri, 22 Aug 2014 15:05:29 +0400 Subject: [PATCH 17/26] naive implementation --- modules/core/src/matmul.cpp | 82 ++++++++++++++++++++++++++++++++- modules/core/src/opencl/gemm.cl | 60 ++++++++++++++++++++++++ 2 files changed, 140 insertions(+), 2 deletions(-) create mode 100644 modules/core/src/opencl/gemm.cl diff --git a/modules/core/src/matmul.cpp b/modules/core/src/matmul.cpp index e3e1720ef0..325d15149e 100644 --- a/modules/core/src/matmul.cpp +++ b/modules/core/src/matmul.cpp @@ -693,7 +693,7 @@ static void GEMMStore_64fc( const Complexd* c_data, size_t c_step, #ifdef HAVE_CLAMDBLAS -static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, +static bool ocl_gemm_amdblas( InputArray matA, InputArray matB, double alpha, InputArray matC, double beta, OutputArray matD, int flags ) { int type = matA.type(), esz = CV_ELEM_SIZE(type); @@ -775,6 +775,79 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, #endif +#ifdef HAVE_OPENCL + +static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, + InputArray matC, double beta, OutputArray matD, int flags ) +{ + int depth = matA.depth(), cn = matA.channels(); + int type = CV_MAKETYPE(depth, cn); + + CV_Assert( type == matB.type() && (type == CV_32FC1 || type == CV_64FC1 || type == CV_32FC2 || type == CV_64FC2) ); + + const ocl::Device & dev = ocl::Device::getDefault(); + bool doubleSupport = dev.doubleFPConfig() > 0; + + if ((!doubleSupport && depth == CV_64F)) + return false; + + bool haveC = matC.kind() != cv::_InputArray::NONE; + Size sizeA = matA.size(), sizeB = matB.size(), sizeC = haveC ? matC.size() : Size(0, 0); + bool atrans = (flags & GEMM_1_T) != 0, btrans = (flags & GEMM_2_T) != 0, ctrans = (flags & GEMM_3_T) != 0; + + if (atrans) + sizeA = Size(sizeA.height, sizeA.width); + if (btrans) + sizeB = Size(sizeB.height, sizeB.width); + if (haveC && ctrans) + sizeC = Size(sizeC.height, sizeC.width); + + Size sizeD(sizeB.width, sizeA.height); + + CV_Assert( matB.type() == type && (!haveC || matC.type() == type) ); + CV_Assert( sizeA.width == sizeB.height && (!haveC || sizeC == sizeD) ); + + String opts = format("-D T=%s -D T1=%s -D cn=%d %s %s", + ocl::typeToStr(type), ocl::typeToStr(depth), cn, + haveC ? "-D HAVE_C" : "", + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); + + ocl::Kernel k("gemm", cv::ocl::core::gemm_oclsrc, opts); + if (k.empty()) + return false; + + matD.create(sizeD, type); + + UMat A = matA.getUMat(), B = matB.getUMat(), D = matD.getUMat(); + + if (atrans) + A = A.t(); + + if (btrans) + B = B.t(); + + if (haveC) + ctrans ? transpose(matC, D) : matC.copyTo(D); + else + D.setTo(Scalar::all(0)); + + if (depth == CV_64F) + k.args(ocl::KernelArg::ReadOnlyNoSize(A), + ocl::KernelArg::ReadOnlyNoSize(B), + ocl::KernelArg::ReadWrite(D), + sizeA.width, alpha, beta); + else + k.args(ocl::KernelArg::ReadOnlyNoSize(A), + ocl::KernelArg::ReadOnlyNoSize(B), + ocl::KernelArg::ReadWrite(D), + sizeA.width, (float)alpha, (float)beta); + + size_t globalsize[2] = { sizeD.width, sizeD.height}; + return k.run(2, globalsize, NULL, false); +} + +#endif + } void cv::gemm( InputArray matA, InputArray matB, double alpha, @@ -783,7 +856,12 @@ void cv::gemm( InputArray matA, InputArray matB, double alpha, #ifdef HAVE_CLAMDBLAS CV_OCL_RUN(ocl::haveAmdBlas() && matA.dims() <= 2 && matB.dims() <= 2 && matC.dims() <= 2 && _matD.isUMat() && matA.cols() > 20 && matA.rows() > 20 && matB.cols() > 20, // since it works incorrect for small sizes - ocl_gemm(matA, matB, alpha, matC, beta, _matD, flags)) + ocl_gemm_amdblas(matA, matB, alpha, matC, beta, _matD, flags)) +#endif + +#ifdef HAVE_OPENCL + CV_OCL_RUN(_matD.isUMat() && matA.dims() <= 2 && matB.dims() <= 2 && matC.dims() <= 2, + ocl_gemm(matA, matB, alpha, matC, beta, _matD, flags)) #endif const int block_lin_size = 128; diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl new file mode 100644 index 0000000000..10e0eea6f3 --- /dev/null +++ b/modules/core/src/opencl/gemm.cl @@ -0,0 +1,60 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. + +// Copyright (C) 2014, Itseez, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. + +#ifdef DOUBLE_SUPPORT +#ifdef cl_amd_fp64 +#pragma OPENCL EXTENSION cl_amd_fp64:enable +#elif defined (cl_khr_fp64) +#pragma OPENCL EXTENSION cl_khr_fp64:enable +#endif +#endif + +#define TSIZE (int)sizeof(T) + +#define IND_A mad24(y, A_step, A_offset) +#define STEP_A 1 + +#define IND_B mad24(x, TSIZE, B_offset) +#define STEP_B B_step / TSIZE + +#if cn==2 +#define MUL(i, a, b)\ + {\ + sum.x += fma(a.x, b.x, - a.y * b.y);\ + sum.y += fma(a.x, b.y, a.y * b.x);\ + } +#else +#define MUL(i, a, b) sum += a * b +#endif + + +__kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, + __global const uchar * B_ptr, int B_step, int B_offset, + __global uchar * D_ptr, int D_step, int D_offset, int D_rows, int D_cols, + int n, T1 alpha, T1 beta) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (x < D_cols && y < D_rows) + { + __global const T* A = (__global const T*)(A_ptr + IND_A); + __global const T* B = (__global const T*)(B_ptr + IND_B); + + T sum = (T)(0); + + for (int i = 0; i < n; ++i) + MUL(i, A[i*STEP_A], B[i*STEP_B]); + + __global T* D = (__global T*)(D_ptr + mad24(y, D_step, mad24(x, TSIZE, D_offset))); +#if HAVE_C + D[0] = mad(alpha, sum, D[0]*beta); +#else + D[0] = alpha * sum; +#endif + } +} From 2d89df18041713c14962da572cc2948123a02efd Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Wed, 27 Aug 2014 14:58:01 +0400 Subject: [PATCH 18/26] use local memory --- modules/core/src/matmul.cpp | 8 +++++--- modules/core/src/opencl/gemm.cl | 36 ++++++++++++++++++++++++--------- modules/ts/src/ocl_test.cpp | 2 +- 3 files changed, 33 insertions(+), 13 deletions(-) diff --git a/modules/core/src/matmul.cpp b/modules/core/src/matmul.cpp index 325d15149e..1db4346efe 100644 --- a/modules/core/src/matmul.cpp +++ b/modules/core/src/matmul.cpp @@ -782,6 +782,7 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, { int depth = matA.depth(), cn = matA.channels(); int type = CV_MAKETYPE(depth, cn); + const int block_size = 16; CV_Assert( type == matB.type() && (type == CV_32FC1 || type == CV_64FC1 || type == CV_32FC2 || type == CV_64FC2) ); @@ -807,8 +808,8 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, CV_Assert( matB.type() == type && (!haveC || matC.type() == type) ); CV_Assert( sizeA.width == sizeB.height && (!haveC || sizeC == sizeD) ); - String opts = format("-D T=%s -D T1=%s -D cn=%d %s %s", - ocl::typeToStr(type), ocl::typeToStr(depth), cn, + String opts = format("-D T=%s -D T1=%s -D cn=%d -D LOCAL_SIZE=%d %s %s", + ocl::typeToStr(type), ocl::typeToStr(depth), cn, block_size, haveC ? "-D HAVE_C" : "", doubleSupport ? " -D DOUBLE_SUPPORT" : ""); @@ -843,7 +844,8 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, sizeA.width, (float)alpha, (float)beta); size_t globalsize[2] = { sizeD.width, sizeD.height}; - return k.run(2, globalsize, NULL, false); + size_t localsize[2] = { block_size, block_size}; + return k.run(2, globalsize, localsize, false); } #endif diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl index 10e0eea6f3..b2437de7e1 100644 --- a/modules/core/src/opencl/gemm.cl +++ b/modules/core/src/opencl/gemm.cl @@ -28,7 +28,7 @@ sum.y += fma(a.x, b.y, a.y * b.x);\ } #else -#define MUL(i, a, b) sum += a * b +#define MUL(i, a, b) sum = fma(a, b, sum); #endif @@ -40,16 +40,34 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, int x = get_global_id(0); int y = get_global_id(1); + int lidx = get_local_id(0); + int lidy = get_local_id(1); + + __global const T* A = (__global const T*)(A_ptr + IND_A); + __global const T* B = (__global const T*)(B_ptr + IND_B); + + T sum = (T)(0); + __local T a_local[LOCAL_SIZE*LOCAL_SIZE]; + __local T b_local[LOCAL_SIZE*LOCAL_SIZE]; + + for (int p = 0; p < (n+LOCAL_SIZE-1)/LOCAL_SIZE; ++p) + { + a_local[mad24(lidy, LOCAL_SIZE, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)]; + b_local[mad24(lidy, LOCAL_SIZE, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B]; + + barrier(CLK_LOCAL_MEM_FENCE); + + if (x < D_cols && y < D_rows) + { + for (int i = 0; i < LOCAL_SIZE && p * LOCAL_SIZE + i < n; ++i) + MUL(i, a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); + } + + barrier(CLK_LOCAL_MEM_FENCE); + } + if (x < D_cols && y < D_rows) { - __global const T* A = (__global const T*)(A_ptr + IND_A); - __global const T* B = (__global const T*)(B_ptr + IND_B); - - T sum = (T)(0); - - for (int i = 0; i < n; ++i) - MUL(i, A[i*STEP_A], B[i*STEP_B]); - __global T* D = (__global T*)(D_ptr + mad24(y, D_step, mad24(x, TSIZE, D_offset))); #if HAVE_C D[0] = mad(alpha, sum, D[0]*beta); diff --git a/modules/ts/src/ocl_test.cpp b/modules/ts/src/ocl_test.cpp index d429d4bc8c..3a3c08cf0f 100644 --- a/modules/ts/src/ocl_test.cpp +++ b/modules/ts/src/ocl_test.cpp @@ -48,7 +48,7 @@ namespace ocl { using namespace cv; -int test_loop_times = 1; // TODO Read from command line / environment +int test_loop_times = 10; // TODO Read from command line / environment #ifdef HAVE_OPENCL From c5a2879ce04582811070c2fd29e364abb9d5dc0a Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Thu, 4 Sep 2014 12:36:23 +0400 Subject: [PATCH 19/26] use vectors --- modules/core/src/matmul.cpp | 40 ++++++++++++--------- modules/core/src/opencl/gemm.cl | 62 +++++++++++++++++++++++++-------- modules/ts/src/ocl_test.cpp | 2 +- 3 files changed, 71 insertions(+), 33 deletions(-) diff --git a/modules/core/src/matmul.cpp b/modules/core/src/matmul.cpp index 1db4346efe..6d2adc8690 100644 --- a/modules/core/src/matmul.cpp +++ b/modules/core/src/matmul.cpp @@ -782,7 +782,6 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, { int depth = matA.depth(), cn = matA.channels(); int type = CV_MAKETYPE(depth, cn); - const int block_size = 16; CV_Assert( type == matB.type() && (type == CV_32FC1 || type == CV_64FC1 || type == CV_32FC2 || type == CV_64FC2) ); @@ -808,14 +807,8 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, CV_Assert( matB.type() == type && (!haveC || matC.type() == type) ); CV_Assert( sizeA.width == sizeB.height && (!haveC || sizeC == sizeD) ); - String opts = format("-D T=%s -D T1=%s -D cn=%d -D LOCAL_SIZE=%d %s %s", - ocl::typeToStr(type), ocl::typeToStr(depth), cn, block_size, - haveC ? "-D HAVE_C" : "", - doubleSupport ? " -D DOUBLE_SUPPORT" : ""); - - ocl::Kernel k("gemm", cv::ocl::core::gemm_oclsrc, opts); - if (k.empty()) - return false; + int max_wg_size = (int)dev.maxWorkGroupSize(); + int block_size = (max_wg_size / (32*cn) < 32) ? (max_wg_size / (16*cn) < 16) ? (max_wg_size / (8*cn) < 8) ? 1 : 8 : 16 : 32; matD.create(sizeD, type); @@ -832,24 +825,37 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, else D.setTo(Scalar::all(0)); + int vectorWidths[] = { 4, 4, 2, 2, 1, 4, cn, -1 }; + + int kercn = ocl::checkOptimalVectorWidth(vectorWidths, B, D); + + String opts = format("-D T=%s -D T1=%s -D WT=%s -D cn=%d -D kercn=%d -D LOCAL_SIZE=%d %s %s %s", + ocl::typeToStr(type), ocl::typeToStr(depth), ocl::typeToStr(CV_MAKETYPE(depth, kercn)), + cn, kercn, block_size, + (sizeA.width % block_size !=0) ? "-D NO_MULT" : "", + haveC ? "-D HAVE_C" : "", + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); + + ocl::Kernel k("gemm", cv::ocl::core::gemm_oclsrc, opts); + if (k.empty()) + return false; + if (depth == CV_64F) k.args(ocl::KernelArg::ReadOnlyNoSize(A), - ocl::KernelArg::ReadOnlyNoSize(B), - ocl::KernelArg::ReadWrite(D), + ocl::KernelArg::ReadOnlyNoSize(B, cn, kercn), + ocl::KernelArg::ReadWrite(D, cn, kercn), sizeA.width, alpha, beta); else k.args(ocl::KernelArg::ReadOnlyNoSize(A), - ocl::KernelArg::ReadOnlyNoSize(B), - ocl::KernelArg::ReadWrite(D), + ocl::KernelArg::ReadOnlyNoSize(B, cn, kercn), + ocl::KernelArg::ReadWrite(D, cn, kercn), sizeA.width, (float)alpha, (float)beta); - size_t globalsize[2] = { sizeD.width, sizeD.height}; + size_t globalsize[2] = { sizeD.width * cn / kercn, sizeD.height}; size_t localsize[2] = { block_size, block_size}; - return k.run(2, globalsize, localsize, false); + return k.run(2, globalsize, block_size!=1 ? localsize : NULL, false); } - #endif - } void cv::gemm( InputArray matA, InputArray matB, double alpha, diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl index b2437de7e1..ddd18adaf2 100644 --- a/modules/core/src/opencl/gemm.cl +++ b/modules/core/src/opencl/gemm.cl @@ -13,21 +13,30 @@ #endif #endif -#define TSIZE (int)sizeof(T) +#define TSIZE (int)sizeof(T) +#define WTSIZE (int)sizeof(WT) #define IND_A mad24(y, A_step, A_offset) -#define STEP_A 1 - -#define IND_B mad24(x, TSIZE, B_offset) -#define STEP_B B_step / TSIZE +#define IND_B mad24(x, WTSIZE, B_offset) +#define STEP_B B_step / WTSIZE #if cn==2 +#if kercn==2 #define MUL(i, a, b)\ {\ sum.x += fma(a.x, b.x, - a.y * b.y);\ sum.y += fma(a.x, b.y, a.y * b.x);\ } #else +#define MUL(i, a, b)\ + {\ + sum.x += fma(a.x, b.x, - a.y * b.y);\ + sum.y += fma(a.x, b.y, a.y * b.x);\ + sum.z += fma(a.x, b.z, - a.y * b.w);\ + sum.w += fma(a.x, b.w, a.y * b.z);\ + } +#endif +#else #define MUL(i, a, b) sum = fma(a, b, sum); #endif @@ -44,22 +53,44 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, int lidy = get_local_id(1); __global const T* A = (__global const T*)(A_ptr + IND_A); - __global const T* B = (__global const T*)(B_ptr + IND_B); + __global const WT* B = (__global const WT*)(B_ptr + IND_B); - T sum = (T)(0); - __local T a_local[LOCAL_SIZE*LOCAL_SIZE]; - __local T b_local[LOCAL_SIZE*LOCAL_SIZE]; + WT sum = (WT)(0); - for (int p = 0; p < (n+LOCAL_SIZE-1)/LOCAL_SIZE; ++p) +#if LOCAL_SIZE == 1 + + if (x < D_cols && y < D_rows) { - a_local[mad24(lidy, LOCAL_SIZE, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)]; - b_local[mad24(lidy, LOCAL_SIZE, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B]; + for (int i = 0; i < n; ++i) + MUL(i, A[i], B[i*STEP_B]); +#else + + __local T a_local[LOCAL_SIZE*LOCAL_SIZE]; + __local WT b_local[LOCAL_SIZE*LOCAL_SIZE]; + + int reps; +#if NO_MULT + reps = (n + LOCAL_SIZE-1)/LOCAL_SIZE; +#else + reps = n/LOCAL_SIZE; +#endif + + for (int p = 0; p < reps; ++p) + { + if (p * LOCAL_SIZE + lidx < n && y < D_rows) + a_local[mad24(lidy, LOCAL_SIZE, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)]; + if (p * LOCAL_SIZE + lidy < n && x < D_cols) + b_local[mad24(lidy, LOCAL_SIZE, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B]; barrier(CLK_LOCAL_MEM_FENCE); if (x < D_cols && y < D_rows) { - for (int i = 0; i < LOCAL_SIZE && p * LOCAL_SIZE + i < n; ++i) + for (int i = 0; i < LOCAL_SIZE +#if NO_MULT + && p * LOCAL_SIZE + i < n +#endif + ; ++i) MUL(i, a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); } @@ -68,11 +99,12 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, if (x < D_cols && y < D_rows) { - __global T* D = (__global T*)(D_ptr + mad24(y, D_step, mad24(x, TSIZE, D_offset))); +#endif + __global WT* D = (__global WT*)(D_ptr + mad24(y, D_step, mad24(x, WTSIZE, D_offset))); #if HAVE_C D[0] = mad(alpha, sum, D[0]*beta); #else D[0] = alpha * sum; #endif } -} +} \ No newline at end of file diff --git a/modules/ts/src/ocl_test.cpp b/modules/ts/src/ocl_test.cpp index 3a3c08cf0f..d429d4bc8c 100644 --- a/modules/ts/src/ocl_test.cpp +++ b/modules/ts/src/ocl_test.cpp @@ -48,7 +48,7 @@ namespace ocl { using namespace cv; -int test_loop_times = 10; // TODO Read from command line / environment +int test_loop_times = 1; // TODO Read from command line / environment #ifdef HAVE_OPENCL From 65b8a1cb3756ba908440efaa4b18eed4037e25c6 Mon Sep 17 00:00:00 2001 From: ElenaGvozdeva Date: Thu, 16 Oct 2014 10:24:44 +0300 Subject: [PATCH 20/26] Some small fixes --- modules/core/src/matmul.cpp | 7 ++----- modules/core/src/opencl/gemm.cl | 18 +++++++++--------- 2 files changed, 11 insertions(+), 14 deletions(-) diff --git a/modules/core/src/matmul.cpp b/modules/core/src/matmul.cpp index 6d2adc8690..fec2ecb8b8 100644 --- a/modules/core/src/matmul.cpp +++ b/modules/core/src/matmul.cpp @@ -788,7 +788,7 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, const ocl::Device & dev = ocl::Device::getDefault(); bool doubleSupport = dev.doubleFPConfig() > 0; - if ((!doubleSupport && depth == CV_64F)) + if (!doubleSupport && depth == CV_64F) return false; bool haveC = matC.kind() != cv::_InputArray::NONE; @@ -804,7 +804,7 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, Size sizeD(sizeB.width, sizeA.height); - CV_Assert( matB.type() == type && (!haveC || matC.type() == type) ); + CV_Assert( !haveC || matC.type() == type ); CV_Assert( sizeA.width == sizeB.height && (!haveC || sizeC == sizeD) ); int max_wg_size = (int)dev.maxWorkGroupSize(); @@ -822,11 +822,8 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, if (haveC) ctrans ? transpose(matC, D) : matC.copyTo(D); - else - D.setTo(Scalar::all(0)); int vectorWidths[] = { 4, 4, 2, 2, 1, 4, cn, -1 }; - int kercn = ocl::checkOptimalVectorWidth(vectorWidths, B, D); String opts = format("-D T=%s -D T1=%s -D WT=%s -D cn=%d -D kercn=%d -D LOCAL_SIZE=%d %s %s %s", diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl index ddd18adaf2..0961628a49 100644 --- a/modules/core/src/opencl/gemm.cl +++ b/modules/core/src/opencl/gemm.cl @@ -22,13 +22,13 @@ #if cn==2 #if kercn==2 -#define MUL(i, a, b)\ +#define MUL(a, b)\ {\ sum.x += fma(a.x, b.x, - a.y * b.y);\ sum.y += fma(a.x, b.y, a.y * b.x);\ } #else -#define MUL(i, a, b)\ +#define MUL(a, b)\ {\ sum.x += fma(a.x, b.x, - a.y * b.y);\ sum.y += fma(a.x, b.y, a.y * b.x);\ @@ -37,7 +37,7 @@ } #endif #else -#define MUL(i, a, b) sum = fma(a, b, sum); +#define MUL(a, b) sum = fma(a, b, sum); #endif @@ -62,7 +62,7 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, if (x < D_cols && y < D_rows) { for (int i = 0; i < n; ++i) - MUL(i, A[i], B[i*STEP_B]); + MUL(A[i], B[i*STEP_B]); #else __local T a_local[LOCAL_SIZE*LOCAL_SIZE]; @@ -86,14 +86,14 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, if (x < D_cols && y < D_rows) { - for (int i = 0; i < LOCAL_SIZE #if NO_MULT - && p * LOCAL_SIZE + i < n + int ie = min(LOCAL_SIZE, n - p * LOCAL_SIZE); + for (int i = 0; i < ie; ++i) +#else + for (int i = 0; i < LOCAL_SIZE; ++i) #endif - ; ++i) - MUL(i, a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); + MUL(a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); } - barrier(CLK_LOCAL_MEM_FENCE); } From 0aab7795326c5a56039c6a7eed7d4bb7f78277e3 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Fri, 24 Oct 2014 11:03:10 +0300 Subject: [PATCH 21/26] Overload PlaneWarper::buildMaps method from base class --- .../stitching/include/opencv2/stitching/detail/warpers.hpp | 1 + modules/stitching/src/opencl/warpers.cl | 4 ++-- modules/stitching/src/warpers.cpp | 5 +++++ 3 files changed, 8 insertions(+), 2 deletions(-) diff --git a/modules/stitching/include/opencv2/stitching/detail/warpers.hpp b/modules/stitching/include/opencv2/stitching/detail/warpers.hpp index c8869f116d..ac9e256ad7 100644 --- a/modules/stitching/include/opencv2/stitching/detail/warpers.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/warpers.hpp @@ -135,6 +135,7 @@ public: Point2f warpPoint(const Point2f &pt, InputArray K, InputArray R, InputArray T); virtual Rect buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray xmap, OutputArray ymap); + Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap); virtual Point warp(InputArray src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, OutputArray dst); diff --git a/modules/stitching/src/opencl/warpers.cl b/modules/stitching/src/opencl/warpers.cl index 7ec87ae2f8..9d245893cf 100644 --- a/modules/stitching/src/opencl/warpers.cl +++ b/modules/stitching/src/opencl/warpers.cl @@ -57,7 +57,7 @@ __kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xm int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset)); float u = tl_u + du; - float x_ = u * scale - ct[0]; + float x_ = fma(u, scale, -ct[0]); float ct1 = 1 - ct[2]; for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step, @@ -67,7 +67,7 @@ __kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xm __global float * ymap = (__global float *)(ymapptr + ymap_index); float v = tl_v + dv; - float y_ = v * scale - ct[1]; + float y_ = fma(v, scale, -ct[1]); float x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * ct1)); float y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * ct1)); diff --git a/modules/stitching/src/warpers.cpp b/modules/stitching/src/warpers.cpp index 2711a012f6..744474ba6e 100644 --- a/modules/stitching/src/warpers.cpp +++ b/modules/stitching/src/warpers.cpp @@ -87,6 +87,11 @@ Point2f PlaneWarper::warpPoint(const Point2f &pt, InputArray K, InputArray R, In return uv; } +Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) +{ + return buildMaps(src_size, K, R, Mat::zeros(3, 1, CV_32FC1), xmap, ymap); +} + Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray _xmap, OutputArray _ymap) { projector_.setCameraParams(K, R, T); From ebfaf4c5d871397d7de10f791907b2e7e0246fbd Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Mon, 27 Oct 2014 15:38:44 +0300 Subject: [PATCH 22/26] Added checking that z is non zero to buildWarpPlaneMaps kernel --- modules/stitching/src/opencl/warpers.cl | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/modules/stitching/src/opencl/warpers.cl b/modules/stitching/src/opencl/warpers.cl index 9d245893cf..a8f04a63cd 100644 --- a/modules/stitching/src/opencl/warpers.cl +++ b/modules/stitching/src/opencl/warpers.cl @@ -73,8 +73,10 @@ __kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xm float y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * ct1)); float z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * ct1)); - x /= z; - y /= z; + if (z != 0) + x /= z, y /= z; + else + x = y = -1; xmap[0] = x; ymap[0] = y; From 1466621f99f0418fb575c68caabb8c200e43904f Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Mon, 27 Oct 2014 14:52:17 +0300 Subject: [PATCH 23/26] Added loading 4 pixels in line instead of 2 to RGB[A] -> YUV(420) kernel --- modules/imgproc/src/color.cpp | 17 +++++-- modules/imgproc/src/opencl/cvtcolor.cl | 62 +++++++++++++++++++++++--- 2 files changed, 71 insertions(+), 8 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index f363189579..dcbfb8f79f 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -4857,6 +4857,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) ocl::Device dev = ocl::Device::getDefault(); int pxPerWIy = dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU) ? 4 : 1; + int pxPerWIx = 1; size_t globalsize[] = { src.cols, (src.rows + pxPerWIy - 1) / pxPerWIy }; cv::String opts = format("-D depth=%d -D scn=%d -D PIX_PER_WI_Y=%d ", @@ -5025,10 +5026,20 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) 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; + _dst.create(dstSz, CV_MAKETYPE(depth, dcn)); + dst = _dst.getUMat(); + + if (dev.isIntel() && src.cols % 4 == 0 && src.step % 4 == 0 && src.offset % 4 == 0 && + dst.step % 4 == 0 && dst.offset % 4 == 0) + { + pxPerWIx = 2; + } + globalsize[0] = dstSz.width / (2 * pxPerWIx); 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; + opts + format("-D dcn=%d -D bidx=%d -D uidx=%d -D PIX_PER_WI_X=%d", dcn, bidx, uidx, pxPerWIx)); + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst)); + return k.run(2, globalsize, NULL, false); } case COLOR_YUV2RGB_UYVY: case COLOR_YUV2BGR_UYVY: case COLOR_YUV2RGBA_UYVY: case COLOR_YUV2BGRA_UYVY: case COLOR_YUV2RGB_YUY2: case COLOR_YUV2BGR_YUY2: case COLOR_YUV2RGB_YVYU: case COLOR_YUV2BGR_YVYU: diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index c3cfd0d592..e660a52204 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -119,6 +119,10 @@ enum #define yidx 0 #endif +#ifndef PIX_PER_WI_X +#define PIX_PER_WI_X 1 +#endif + #define __CAT(x, y) x##y #define CAT(x, y) __CAT(x, y) @@ -454,7 +458,7 @@ __kernel void RGB2YUV_YV12_IYUV(__global const uchar* srcptr, int src_step, int __global uchar* dstptr, int dst_step, int dst_offset, int rows, int cols) { - int x = get_global_id(0); + int x = get_global_id(0) * PIX_PER_WI_X; int y = get_global_id(1) * PIX_PER_WI_Y; if (x < cols/2) @@ -463,6 +467,7 @@ __kernel void RGB2YUV_YV12_IYUV(__global const uchar* srcptr, int src_step, int 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)}; + __constant float* coeffs = c_RGB2YUVCoeffs_420; #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) @@ -477,12 +482,61 @@ __kernel void RGB2YUV_YV12_IYUV(__global const uchar* srcptr, int src_step, int __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); +#if PIX_PER_WI_X == 2 + int s11 = *((__global const int*) src1); + int s12 = *((__global const int*) src1 + 1); + int s13 = *((__global const int*) src1 + 2); +#if scn == 4 + int s14 = *((__global const int*) src1 + 3); +#endif + int s21 = *((__global const int*) src2); + int s22 = *((__global const int*) src2 + 1); + int s23 = *((__global const int*) src2 + 2); +#if scn == 4 + int s24 = *((__global const int*) src2 + 3); +#endif + float src_pix1[scn * 4], src_pix2[scn * 4]; + + *((float4*) src_pix1) = convert_float4(as_uchar4(s11)); + *((float4*) src_pix1 + 1) = convert_float4(as_uchar4(s12)); + *((float4*) src_pix1 + 2) = convert_float4(as_uchar4(s13)); +#if scn == 4 + *((float4*) src_pix1 + 3) = convert_float4(as_uchar4(s14)); +#endif + *((float4*) src_pix2) = convert_float4(as_uchar4(s21)); + *((float4*) src_pix2 + 1) = convert_float4(as_uchar4(s22)); + *((float4*) src_pix2 + 2) = convert_float4(as_uchar4(s23)); +#if scn == 4 + *((float4*) src_pix2 + 3) = convert_float4(as_uchar4(s24)); +#endif + uchar4 y1, y2; + y1.x = convert_uchar_sat(fma(coeffs[0], src_pix1[ 2-bidx], fma(coeffs[1], src_pix1[ 1], fma(coeffs[2], src_pix1[ bidx], 16.5f)))); + y1.y = convert_uchar_sat(fma(coeffs[0], src_pix1[ scn+2-bidx], fma(coeffs[1], src_pix1[ scn+1], fma(coeffs[2], src_pix1[ scn+bidx], 16.5f)))); + y1.z = convert_uchar_sat(fma(coeffs[0], src_pix1[2*scn+2-bidx], fma(coeffs[1], src_pix1[2*scn+1], fma(coeffs[2], src_pix1[2*scn+bidx], 16.5f)))); + y1.w = convert_uchar_sat(fma(coeffs[0], src_pix1[3*scn+2-bidx], fma(coeffs[1], src_pix1[3*scn+1], fma(coeffs[2], src_pix1[3*scn+bidx], 16.5f)))); + y2.x = convert_uchar_sat(fma(coeffs[0], src_pix2[ 2-bidx], fma(coeffs[1], src_pix2[ 1], fma(coeffs[2], src_pix2[ bidx], 16.5f)))); + y2.y = convert_uchar_sat(fma(coeffs[0], src_pix2[ scn+2-bidx], fma(coeffs[1], src_pix2[ scn+1], fma(coeffs[2], src_pix2[ scn+bidx], 16.5f)))); + y2.z = convert_uchar_sat(fma(coeffs[0], src_pix2[2*scn+2-bidx], fma(coeffs[1], src_pix2[2*scn+1], fma(coeffs[2], src_pix2[2*scn+bidx], 16.5f)))); + y2.w = convert_uchar_sat(fma(coeffs[0], src_pix2[3*scn+2-bidx], fma(coeffs[1], src_pix2[3*scn+1], fma(coeffs[2], src_pix2[3*scn+bidx], 16.5f)))); + + *((__global int*) ydst1) = as_int(y1); + *((__global int*) ydst2) = as_int(y2); + + float uv[4] = { fma(coeffs[3], src_pix1[ 2-bidx], fma(coeffs[4], src_pix1[ 1], fma(coeffs[5], src_pix1[ bidx], 128.5f))), + fma(coeffs[5], src_pix1[ 2-bidx], fma(coeffs[6], src_pix1[ 1], fma(coeffs[7], src_pix1[ bidx], 128.5f))), + fma(coeffs[3], src_pix1[2*scn+2-bidx], fma(coeffs[4], src_pix1[2*scn+1], fma(coeffs[5], src_pix1[2*scn+bidx], 128.5f))), + fma(coeffs[5], src_pix1[2*scn+2-bidx], fma(coeffs[6], src_pix1[2*scn+1], fma(coeffs[7], src_pix1[2*scn+bidx], 128.5f))) }; + + udst[0] = convert_uchar_sat(uv[uidx] ); + vdst[0] = convert_uchar_sat(uv[1 - uidx]); + udst[1] = convert_uchar_sat(uv[2 + uidx]); + vdst[1] = convert_uchar_sat(uv[3 - uidx]); +#else float4 src_pix1 = convert_float4(vload4(0, src1)); float4 src_pix2 = convert_float4(vload4(0, src1+scn)); float4 src_pix3 = convert_float4(vload4(0, src2)); float4 src_pix4 = convert_float4(vload4(0, src2+scn)); - __constant float* coeffs = c_RGB2YUVCoeffs_420; ydst1[0] = convert_uchar_sat(fma(coeffs[0], src_pix1.R_COMP, fma(coeffs[1], src_pix1.G_COMP, fma(coeffs[2], src_pix1.B_COMP, 16.5f)))); ydst1[1] = convert_uchar_sat(fma(coeffs[0], src_pix2.R_COMP, fma(coeffs[1], src_pix2.G_COMP, fma(coeffs[2], src_pix2.B_COMP, 16.5f)))); ydst2[0] = convert_uchar_sat(fma(coeffs[0], src_pix3.R_COMP, fma(coeffs[1], src_pix3.G_COMP, fma(coeffs[2], src_pix3.B_COMP, 16.5f)))); @@ -493,7 +547,7 @@ __kernel void RGB2YUV_YV12_IYUV(__global const uchar* srcptr, int src_step, int udst[0] = convert_uchar_sat(uv[uidx] ); vdst[0] = convert_uchar_sat(uv[1-uidx]); - +#endif ++y; src_index += 2*src_step; ydst_index += 2*dst_step; @@ -522,7 +576,6 @@ __kernel void YUV2RGB_422(__global const uchar* srcptr, int src_step, int src_of float U = ((float) src[uidx]) - HALF_MAX; float V = ((float) src[(2 + uidx) % 4]) - HALF_MAX; - __constant float* coeffs = c_YUV2RGBCoeffs_420; float ruv = fma(coeffs[4], V, 0.5f); float guv = fma(coeffs[3], V, fma(coeffs[2], U, 0.5f)); @@ -535,7 +588,6 @@ __kernel void YUV2RGB_422(__global const uchar* srcptr, int src_step, int src_of #if dcn == 4 dst[3] = 255; #endif - float y01 = max(0.f, ((float) src[yidx + 2]) - 16.f) * coeffs[0]; dst[dcn + 2 - bidx] = convert_uchar_sat(y01 + ruv); dst[dcn + 1] = convert_uchar_sat(y01 + guv); From 643c906f3d79933480fcbebdbf16fc431bece1da Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Tue, 28 Oct 2014 15:07:51 +0300 Subject: [PATCH 24/26] Added optimized loading to YUV2RGB_422 kernel --- modules/imgproc/src/color.cpp | 3 ++- modules/imgproc/src/opencl/cvtcolor.cl | 17 ++++++++++++++--- 2 files changed, 16 insertions(+), 4 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index dcbfb8f79f..f0a8fd8584 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -5060,7 +5060,8 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) CV_Assert( scn == 2 && depth == CV_8U ); k.create("YUV2RGB_422", ocl::imgproc::cvtcolor_oclsrc, - opts + format("-D dcn=%d -D bidx=%d -D uidx=%d -D yidx=%d", dcn, bidx, uidx, yidx)); + opts + format("-D dcn=%d -D bidx=%d -D uidx=%d -D yidx=%d%s", dcn, bidx, uidx, yidx, + src.offset % 4 == 0 && src.step % 4 == 0 ? " -D USE_OPTIMIZED_LOAD" : "")); break; } case COLOR_BGR2YCrCb: diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index e660a52204..a7cc776503 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -573,22 +573,33 @@ __kernel void YUV2RGB_422(__global const uchar* srcptr, int src_step, int src_of { if (y < rows ) { + __constant float* coeffs = c_YUV2RGBCoeffs_420; + +#ifndef USE_OPTIMIZED_LOAD float U = ((float) src[uidx]) - HALF_MAX; float V = ((float) src[(2 + uidx) % 4]) - HALF_MAX; + float y00 = max(0.f, ((float) src[yidx]) - 16.f) * coeffs[0]; + float y01 = max(0.f, ((float) src[yidx + 2]) - 16.f) * coeffs[0]; +#else + int load_src = *((__global int*) src); + float vec_src[4] = { load_src & 0xff, (load_src >> 8) & 0xff, (load_src >> 16) & 0xff, (load_src >> 24) & 0xff}; + float U = vec_src[uidx] - HALF_MAX; + float V = vec_src[(2 + uidx) % 4] - HALF_MAX; + float y00 = max(0.f, vec_src[yidx] - 16.f) * coeffs[0]; + float y01 = max(0.f, vec_src[yidx + 2] - 16.f) * coeffs[0]; +#endif - __constant float* coeffs = c_YUV2RGBCoeffs_420; float ruv = fma(coeffs[4], V, 0.5f); float guv = fma(coeffs[3], V, fma(coeffs[2], U, 0.5f)); float buv = fma(coeffs[1], U, 0.5f); - float y00 = max(0.f, ((float) src[yidx]) - 16.f) * coeffs[0]; dst[2 - bidx] = convert_uchar_sat(y00 + ruv); dst[1] = convert_uchar_sat(y00 + guv); dst[bidx] = convert_uchar_sat(y00 + buv); #if dcn == 4 dst[3] = 255; #endif - float y01 = max(0.f, ((float) src[yidx + 2]) - 16.f) * coeffs[0]; + dst[dcn + 2 - bidx] = convert_uchar_sat(y01 + ruv); dst[dcn + 1] = convert_uchar_sat(y01 + guv); dst[dcn + bidx] = convert_uchar_sat(y01 + buv); From d88fdd0378c280763172652931b8d06abe3ebbb5 Mon Sep 17 00:00:00 2001 From: ElenaGvozdeva Date: Tue, 28 Oct 2014 15:18:31 +0300 Subject: [PATCH 25/26] use LOCAL_SIZE+1 --- modules/core/src/opencl/gemm.cl | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl index 0961628a49..fc050547be 100644 --- a/modules/core/src/opencl/gemm.cl +++ b/modules/core/src/opencl/gemm.cl @@ -20,6 +20,8 @@ #define IND_B mad24(x, WTSIZE, B_offset) #define STEP_B B_step / WTSIZE +#define LOCAL_SIZE_ODD (LOCAL_SIZE + 1) + #if cn==2 #if kercn==2 #define MUL(a, b)\ @@ -65,8 +67,8 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, MUL(A[i], B[i*STEP_B]); #else - __local T a_local[LOCAL_SIZE*LOCAL_SIZE]; - __local WT b_local[LOCAL_SIZE*LOCAL_SIZE]; + __local T a_local[LOCAL_SIZE_ODD*LOCAL_SIZE]; + __local WT b_local[LOCAL_SIZE_ODD*LOCAL_SIZE]; int reps; #if NO_MULT @@ -78,9 +80,9 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, for (int p = 0; p < reps; ++p) { if (p * LOCAL_SIZE + lidx < n && y < D_rows) - a_local[mad24(lidy, LOCAL_SIZE, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)]; + a_local[mad24(lidy, LOCAL_SIZE_ODD, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)]; if (p * LOCAL_SIZE + lidy < n && x < D_cols) - b_local[mad24(lidy, LOCAL_SIZE, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B]; + b_local[mad24(lidy, LOCAL_SIZE_ODD, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B]; barrier(CLK_LOCAL_MEM_FENCE); @@ -92,7 +94,7 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, #else for (int i = 0; i < LOCAL_SIZE; ++i) #endif - MUL(a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); + MUL(a_local[mad24(lidy, LOCAL_SIZE_ODD, i)], b_local[mad24(i, LOCAL_SIZE_ODD, lidx)]); } barrier(CLK_LOCAL_MEM_FENCE); } From 967a88759fde8d3a5226511d2772e6e964177f09 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Tue, 28 Oct 2014 16:06:36 +0300 Subject: [PATCH 26/26] Don't process empty src image in GFTT --- modules/imgproc/src/featureselect.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/modules/imgproc/src/featureselect.cpp b/modules/imgproc/src/featureselect.cpp index 4a234f3a96..e51859e012 100644 --- a/modules/imgproc/src/featureselect.cpp +++ b/modules/imgproc/src/featureselect.cpp @@ -275,6 +275,12 @@ void cv::goodFeaturesToTrack( InputArray _image, OutputArray _corners, _mask, blockSize, useHarrisDetector, harrisK)) Mat image = _image.getMat(), eig, tmp; + if (image.empty()) + { + _corners.release(); + return; + } + if( useHarrisDetector ) cornerHarris( image, eig, blockSize, 3, harrisK ); else