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) \ { \