mirror of
https://github.com/opencv/opencv.git
synced 2025-01-18 06:03:15 +08:00
Merge pull request #1886 from ilya-lavrenov:tapi_lut
This commit is contained in:
commit
b27d7d2a61
@ -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()];
|
||||
|
71
modules/core/src/opencl/lut.cl
Normal file
71
modules/core/src/opencl/lut.cl
Normal file
@ -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
|
||||
}
|
||||
}
|
@ -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()));
|
||||
|
||||
|
@ -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()
|
||||
{
|
||||
|
@ -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()
|
||||
{
|
||||
|
@ -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) \
|
||||
{ \
|
||||
|
Loading…
Reference in New Issue
Block a user