From 4eba74ee19ccb570bf7fd1eaaa016c4dcd113ca9 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Thu, 28 Nov 2013 00:57:36 +0400
Subject: [PATCH] added cv::LUT to T-API
---
modules/core/src/convert.cpp | 42 ++++++++++--
modules/core/src/opencl/lut.cl | 71 +++++++++++++++++++++
modules/core/test/ocl/test_arithm.cpp | 74 ++++++++++++++++++++--
modules/imgproc/test/ocl/test_color.cpp | 4 +-
modules/imgproc/test/ocl/test_warp.cpp | 4 +-
modules/ts/include/opencv2/ts/ocl_test.hpp | 4 +-
6 files changed, 182 insertions(+), 17 deletions(-)
create mode 100644 modules/core/src/opencl/lut.cl
diff --git a/modules/core/src/convert.cpp b/modules/core/src/convert.cpp
index b6c413040..af7e042e1 100644
--- a/modules/core/src/convert.cpp
+++ b/modules/core/src/convert.cpp
@@ -41,6 +41,7 @@
//M*/
#include "precomp.hpp"
+#include "opencl_kernels.hpp"
namespace cv
{
@@ -1209,16 +1210,45 @@ static LUTFunc lutTab[] =
}
+namespace cv {
+
+static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst)
+{
+ int dcn = _dst.channels(), lcn = _lut.channels(), dtype = _dst.type();
+
+ if (_src.dims() > 2)
+ return false;
+
+ UMat src = _src.getUMat(), lut = _lut.getUMat();
+ _dst.create(src.size(), dtype);
+ UMat dst = _dst.getUMat();
+
+ ocl::Kernel k("LUT", ocl::core::lut_oclsrc,
+ format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", dcn, lcn,
+ ocl::typeToStr(src.depth()), ocl::typeToStr(dst.depth())));
+ k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut),
+ ocl::KernelArg::WriteOnly(dst));
+
+ size_t globalSize[2] = { dst.cols, dst.rows };
+ return k.run(2, globalSize, NULL, false);
+}
+
+} // cv
+
void cv::LUT( InputArray _src, InputArray _lut, OutputArray _dst )
{
- Mat src = _src.getMat(), lut = _lut.getMat();
- int cn = src.channels();
- int lutcn = lut.channels();
+ int cn = _src.channels(), depth = _src.depth();
+ int lutcn = _lut.channels();
CV_Assert( (lutcn == cn || lutcn == 1) &&
- lut.total() == 256 && lut.isContinuous() &&
- (src.depth() == CV_8U || src.depth() == CV_8S) );
- _dst.create( src.dims, src.size, CV_MAKETYPE(lut.depth(), cn));
+ _lut.total() == 256 && _lut.isContinuous() &&
+ (depth == CV_8U || depth == CV_8S) );
+
+ if (ocl::useOpenCL() && _dst.isUMat() && ocl_LUT(_src, _lut, _dst))
+ return;
+
+ Mat src = _src.getMat(), lut = _lut.getMat();
+ _dst.create(src.dims, src.size, CV_MAKETYPE(_lut.depth(), cn));
Mat dst = _dst.getMat();
LUTFunc func = lutTab[lut.depth()];
diff --git a/modules/core/src/opencl/lut.cl b/modules/core/src/opencl/lut.cl
new file mode 100644
index 000000000..8d5874210
--- /dev/null
+++ b/modules/core/src/opencl/lut.cl
@@ -0,0 +1,71 @@
+// License Agreement
+// For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+// Niko Li, newlife20080214@gmail.com
+// Rock Li, Rock.li@amd.com
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+// * Redistribution's of source code must retain the above copyright notice,
+// this list of conditions and the following disclaimer.
+//
+// * Redistribution's in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// * The name of the copyright holders may not be used to endorse or promote products
+// derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//
+
+#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
+
+__kernel void LUT(__global const uchar * srcptr, int src_step, int src_offset,
+ __global const uchar * lutptr, int lut_step, int lut_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);
+
+ if (x < cols && y < rows)
+ {
+ int src_index = mad24(y, src_step, src_offset + x * (int)sizeof(srcT) * dcn);
+ int dst_index = mad24(y, dst_step, dst_offset + x * (int)sizeof(dstT) * dcn);
+
+ __global const srcT * src = (__global const srcT *)(srcptr + src_index);
+ __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);
+ __global dstT * dst = (__global dstT *)(dstptr + dst_index);
+
+#if lcn == 1
+ #pragma unroll
+ for (int cn = 0; cn < dcn; ++cn)
+ dst[cn] = lut[src[cn]];
+#else
+ #pragma unroll
+ for (int cn = 0; cn < dcn; ++cn)
+ dst[cn] = lut[src[cn] * dcn + cn];
+#endif
+ }
+}
diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp
index 51f50c1c4..1194bd0f6 100644
--- a/modules/core/test/ocl/test_arithm.cpp
+++ b/modules/core/test/ocl/test_arithm.cpp
@@ -47,6 +47,69 @@
namespace cvtest {
namespace ocl {
+//////////////////////////////// LUT /////////////////////////////////////////////////
+
+PARAM_TEST_CASE(Lut, MatDepth, MatDepth, Channels, bool, bool)
+{
+ int src_depth, lut_depth;
+ int cn;
+ bool use_roi, same_cn;
+
+ TEST_DECLARE_INPUT_PARAMETER(src)
+ TEST_DECLARE_INPUT_PARAMETER(lut)
+ TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+ virtual void SetUp()
+ {
+ src_depth = GET_PARAM(0);
+ lut_depth = GET_PARAM(1);
+ cn = GET_PARAM(2);
+ same_cn = GET_PARAM(3);
+ use_roi = GET_PARAM(4);
+ }
+
+ void generateTestData()
+ {
+ const int src_type = CV_MAKE_TYPE(src_depth, cn);
+ const int lut_type = CV_MAKE_TYPE(lut_depth, same_cn ? cn : 1);
+ const int dst_type = CV_MAKE_TYPE(lut_depth, cn);
+
+ Size roiSize = randomSize(1, MAX_VALUE);
+ Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(src, src_roi, roiSize, srcBorder, src_type, 0, 256);
+
+ Size lutRoiSize = Size(256, 1);
+ Border lutBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(lut, lut_roi, lutRoiSize, lutBorder, lut_type, 5, 16);
+
+ Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(dst, dst_roi, roiSize, dstBorder, dst_type, 5, 16);
+
+ UMAT_UPLOAD_INPUT_PARAMETER(src)
+ UMAT_UPLOAD_INPUT_PARAMETER(lut)
+ UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
+ }
+
+ void Near(double threshold = 0.)
+ {
+ EXPECT_MAT_NEAR(dst, udst, threshold);
+ EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold);
+ }
+};
+
+OCL_TEST_P(Lut, Mat)
+{
+ for (int j = 0; j < test_loop_times; j++)
+ {
+ generateTestData();
+
+ OCL_OFF(cv::LUT(src_roi, lut_roi, dst_roi));
+ OCL_ON(cv::LUT(usrc_roi, ulut_roi, udst_roi));
+
+ Near();
+ }
+}
+
///////////////////////// ArithmTestBase ///////////////////////////
PARAM_TEST_CASE(ArithmTestBase, MatDepth, Channels, bool)
@@ -57,11 +120,11 @@ PARAM_TEST_CASE(ArithmTestBase, MatDepth, Channels, bool)
cv::Scalar val;
// declare Mat + UMat mirrors
- TEST_DECLARE_INPUT_PARATEMER(src1)
- TEST_DECLARE_INPUT_PARATEMER(src2)
- TEST_DECLARE_INPUT_PARATEMER(mask)
- TEST_DECLARE_OUTPUT_PARATEMER(dst1)
- TEST_DECLARE_OUTPUT_PARATEMER(dst2)
+ TEST_DECLARE_INPUT_PARAMETER(src1)
+ TEST_DECLARE_INPUT_PARAMETER(src2)
+ TEST_DECLARE_INPUT_PARAMETER(mask)
+ TEST_DECLARE_OUTPUT_PARAMETER(dst1)
+ TEST_DECLARE_OUTPUT_PARAMETER(dst2)
virtual void SetUp()
{
@@ -220,6 +283,7 @@ OCL_TEST_P(Subtract, Scalar_Mask)
//////////////////////////////////////// Instantiation /////////////////////////////////////////
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, ::testing::Values(1, 2, 3, 4), Bool(), Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Add, Combine(OCL_ALL_DEPTHS, ::testing::Values(1, 2, 4), Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Subtract, Combine(OCL_ALL_DEPTHS, ::testing::Values(1, 2, 4), Bool()));
diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp
index cef0c9239..82bd0fc75 100644
--- a/modules/imgproc/test/ocl/test_color.cpp
+++ b/modules/imgproc/test/ocl/test_color.cpp
@@ -60,8 +60,8 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool)
int depth;
bool use_roi;
- TEST_DECLARE_INPUT_PARATEMER(src)
- TEST_DECLARE_OUTPUT_PARATEMER(dst)
+ TEST_DECLARE_INPUT_PARAMETER(src)
+ TEST_DECLARE_OUTPUT_PARAMETER(dst)
virtual void SetUp()
{
diff --git a/modules/imgproc/test/ocl/test_warp.cpp b/modules/imgproc/test/ocl/test_warp.cpp
index 80a55dc45..1bc8f2bb6 100644
--- a/modules/imgproc/test/ocl/test_warp.cpp
+++ b/modules/imgproc/test/ocl/test_warp.cpp
@@ -69,8 +69,8 @@ PARAM_TEST_CASE(Resize, MatType, double, double, Interpolation, bool)
double fx, fy;
bool useRoi;
- TEST_DECLARE_INPUT_PARATEMER(src)
- TEST_DECLARE_OUTPUT_PARATEMER(dst)
+ TEST_DECLARE_INPUT_PARAMETER(src)
+ TEST_DECLARE_OUTPUT_PARAMETER(dst)
virtual void SetUp()
{
diff --git a/modules/ts/include/opencv2/ts/ocl_test.hpp b/modules/ts/include/opencv2/ts/ocl_test.hpp
index 008ce37f6..26dd3df86 100644
--- a/modules/ts/include/opencv2/ts/ocl_test.hpp
+++ b/modules/ts/include/opencv2/ts/ocl_test.hpp
@@ -258,8 +258,8 @@ struct TestUtils
}
};
-#define TEST_DECLARE_INPUT_PARATEMER(name) Mat name, name ## _roi; UMat u ## name, u ## name ## _roi;
-#define TEST_DECLARE_OUTPUT_PARATEMER(name) TEST_DECLARE_INPUT_PARATEMER(name)
+#define TEST_DECLARE_INPUT_PARAMETER(name) Mat name, name ## _roi; UMat u ## name, u ## name ## _roi;
+#define TEST_DECLARE_OUTPUT_PARAMETER(name) TEST_DECLARE_INPUT_PARAMETER(name)
#define UMAT_UPLOAD_INPUT_PARAMETER(name) \
{ \