parent
5c72c74b09
commit
c9e6ed7a29
4 changed files with 365 additions and 0 deletions
@ -0,0 +1,145 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Nathan, liujun@multicorewareinc.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.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#include "precomp.hpp" |
||||
#include "opencl_kernels.hpp" |
||||
|
||||
namespace cv { |
||||
|
||||
template <typename T> |
||||
class BlendLinearInvoker : |
||||
public ParallelLoopBody |
||||
{ |
||||
public: |
||||
BlendLinearInvoker(const Mat & _src1, const Mat & _src2, const Mat & _weights1, |
||||
const Mat & _weights2, Mat & _dst) : |
||||
src1(&_src1), src2(&_src2), weights1(&_weights1), weights2(&_weights2), dst(&_dst) |
||||
{ |
||||
} |
||||
|
||||
virtual void operator() (const Range & range) const |
||||
{ |
||||
int cn = src1->channels(), width = src1->cols * cn; |
||||
|
||||
for (int y = range.start; y < range.end; ++y) |
||||
{ |
||||
const float * const weights1_row = weights1->ptr<float>(y); |
||||
const float * const weights2_row = weights2->ptr<float>(y); |
||||
const T * const src1_row = src1->ptr<T>(y); |
||||
const T * const src2_row = src2->ptr<T>(y); |
||||
T * const dst_row = dst->ptr<T>(y); |
||||
|
||||
for (int x = 0; x < width; ++x) |
||||
{ |
||||
int x1 = x / cn; |
||||
float w1 = weights1_row[x1], w2 = weights2_row[x1]; |
||||
float den = (w1 + w2 + 1e-5f); |
||||
float num = (src1_row[x] * w1 + src2_row[x] * w2); |
||||
|
||||
dst_row[x] = saturate_cast<T>(num / den); |
||||
} |
||||
} |
||||
} |
||||
|
||||
private: |
||||
const BlendLinearInvoker & operator= (const BlendLinearInvoker &); |
||||
BlendLinearInvoker(const BlendLinearInvoker &); |
||||
|
||||
const Mat * src1, * src2, * weights1, * weights2; |
||||
Mat * dst; |
||||
}; |
||||
|
||||
static bool ocl_blendLinear( InputArray _src1, InputArray _src2, InputArray _weights1, InputArray _weights2, OutputArray _dst ) |
||||
{ |
||||
int type = _src1.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); |
||||
|
||||
char cvt[30]; |
||||
ocl::Kernel k("blendLinear", ocl::imgproc::blend_linear_oclsrc, |
||||
format("-D T=%s -D cn=%d -D convertToT=%s", ocl::typeToStr(depth), |
||||
cn, ocl::convertTypeStr(CV_32F, depth, 1, cvt))); |
||||
if (k.empty()) |
||||
return false; |
||||
|
||||
UMat src1 = _src1.getUMat(), src2 = _src2.getUMat(), weights1 = _weights1.getUMat(), |
||||
weights2 = _weights2.getUMat(), dst = _dst.getUMat(); |
||||
|
||||
k.args(ocl::KernelArg::ReadOnlyNoSize(src1), ocl::KernelArg::ReadOnlyNoSize(src2), |
||||
ocl::KernelArg::ReadOnlyNoSize(weights1), ocl::KernelArg::ReadOnlyNoSize(weights2), |
||||
ocl::KernelArg::WriteOnly(dst)); |
||||
|
||||
size_t globalsize[2] = { dst.cols, dst.rows }; |
||||
return k.run(2, globalsize, NULL, false); |
||||
} |
||||
|
||||
} |
||||
|
||||
void cv::blendLinear( InputArray _src1, InputArray _src2, InputArray _weights1, InputArray _weights2, OutputArray _dst ) |
||||
{ |
||||
int type = _src1.type(), depth = CV_MAT_DEPTH(type); |
||||
Size size = _src1.size(); |
||||
|
||||
CV_Assert(depth == CV_8U || depth == CV_32F); |
||||
CV_Assert(size == _src2.size() && size == _weights1.size() && size == _weights2.size()); |
||||
CV_Assert(type == _src2.type() && _weights1.type() == CV_32FC1 && _weights2.type() == CV_32FC1); |
||||
|
||||
_dst.create(size, type); |
||||
|
||||
if (ocl::useOpenCL() && _dst.isUMat() && ocl_blendLinear(_src1, _src2, _weights1, _weights2, _dst)) |
||||
return; |
||||
|
||||
Mat src1 = _src1.getMat(), src2 = _src2.getMat(), weights1 = _weights1.getMat(), |
||||
weights2 = _weights2.getMat(), dst = _dst.getMat(); |
||||
|
||||
if (depth == CV_8U) |
||||
{ |
||||
BlendLinearInvoker<uchar> invoker(src1, src2, weights1, weights2, dst); |
||||
parallel_for_(Range(0, src1.rows), invoker, dst.total()/(double)(1<<16)); |
||||
} |
||||
else if (depth == CV_32F) |
||||
{ |
||||
BlendLinearInvoker<float> invoker(src1, src2, weights1, weights2, dst); |
||||
parallel_for_(Range(0, src1.rows), invoker, dst.total()/(double)(1<<16)); |
||||
} |
||||
} |
@ -0,0 +1,88 @@ |
||||
/*M/////////////////////////////////////////////////////////////////////////////////////// |
||||
// |
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. |
||||
// |
||||
// By downloading, copying, installing or using the software you agree to this license. |
||||
// If you do not agree to this license, do not download, install, |
||||
// copy or use the software. |
||||
// |
||||
// |
||||
// License Agreement |
||||
// For Open Source Computer Vision Library |
||||
// |
||||
// Copyright (C) 2010-2012, MulticoreWare Inc., all rights reserved. |
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. |
||||
// Third party copyrights are property of their respective owners. |
||||
// |
||||
// @Authors |
||||
// Liu Liujun, liujun@multicorewareinc.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. |
||||
// |
||||
//M*/ |
||||
|
||||
#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 noconvert |
||||
|
||||
__kernel void blendLinear(__global const uchar * src1ptr, int src1_step, int src1_offset, |
||||
__global const uchar * src2ptr, int src2_step, int src2_offset, |
||||
__global const uchar * weight1, int weight1_step, int weight1_offset, |
||||
__global const uchar * weight2, int weight2_step, int weight2_offset, |
||||
__global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < dst_cols && y < dst_rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, src1_offset + x * cn * (int)sizeof(T)); |
||||
int src2_index = mad24(y, src2_step, src2_offset + x * cn * (int)sizeof(T)); |
||||
int weight1_index = mad24(y, weight1_step, weight1_offset + x * (int)sizeof(float)); |
||||
int weight2_index = mad24(y, weight2_step, weight2_offset + x * (int)sizeof(float)); |
||||
int dst_index = mad24(y, dst_step, dst_offset + x * cn * (int)sizeof(T)); |
||||
|
||||
float w1 = *(__global const float *)(weight1 + weight1_index), |
||||
w2 = *(__global const float *)(weight2 + weight2_index); |
||||
float den = w1 + w2 + 1e-5f; |
||||
|
||||
__global const T * src1 = (__global const T *)(src1ptr + src1_index); |
||||
__global const T * src2 = (__global const T *)(src2ptr + src2_index); |
||||
__global T * dst = (__global T *)(dstptr + dst_index); |
||||
|
||||
#pragma unroll |
||||
for (int i = 0; i < cn; ++i) |
||||
{ |
||||
float num = w1 * convert_float(src1[i]) + w2 * convert_float(src2[i]); |
||||
dst[i] = convertToT(num / den); |
||||
} |
||||
} |
||||
} |
@ -0,0 +1,129 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Nathan, liujun@multicorewareinc.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.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#include "test_precomp.hpp" |
||||
#include "cvconfig.h" |
||||
#include "opencv2/ts/ocl_test.hpp" |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
namespace cvtest { |
||||
namespace ocl { |
||||
|
||||
PARAM_TEST_CASE(BlendLinear, MatDepth, Channels, bool) |
||||
{ |
||||
int depth, channels; |
||||
bool useRoi; |
||||
|
||||
TEST_DECLARE_INPUT_PARAMETER(src1) |
||||
TEST_DECLARE_INPUT_PARAMETER(src2) |
||||
TEST_DECLARE_INPUT_PARAMETER(weights2) |
||||
TEST_DECLARE_INPUT_PARAMETER(weights1) |
||||
TEST_DECLARE_OUTPUT_PARAMETER(dst) |
||||
|
||||
virtual void SetUp() |
||||
{ |
||||
depth = GET_PARAM(0); |
||||
channels = GET_PARAM(1); |
||||
useRoi = GET_PARAM(2); |
||||
} |
||||
|
||||
void random_roi() |
||||
{ |
||||
const int type = CV_MAKE_TYPE(depth, channels); |
||||
const double upValue = 256; |
||||
|
||||
Size roiSize = randomSize(1, 20); |
||||
Border src1Border = randomBorder(0, useRoi ? MAX_VALUE : 0); |
||||
randomSubMat(src1, src1_roi, roiSize, src1Border, type, -upValue, upValue); |
||||
|
||||
Border src2Border = randomBorder(0, useRoi ? MAX_VALUE : 0); |
||||
randomSubMat(src2, src2_roi, roiSize, src2Border, type, -upValue, upValue); |
||||
|
||||
Border weights1Border = randomBorder(0, useRoi ? MAX_VALUE : 0); |
||||
randomSubMat(weights1, weights1_roi, roiSize, weights1Border, CV_32FC1, -upValue, upValue); |
||||
|
||||
Border weights2Border = randomBorder(0, useRoi ? MAX_VALUE : 0); |
||||
randomSubMat(weights2, weights2_roi, roiSize, weights2Border, CV_32FC1, 1e-2, upValue); |
||||
|
||||
weights2_roi -= weights1_roi; |
||||
CV_Assert(checkNorm(weights2_roi, weights2(Rect(weights2Border.lef, weights2Border.top, |
||||
roiSize.width, roiSize.height))) < 1e-6); |
||||
|
||||
Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); |
||||
randomSubMat(dst, dst_roi, roiSize, dstBorder, type, 5, 16); |
||||
|
||||
UMAT_UPLOAD_INPUT_PARAMETER(src1) |
||||
UMAT_UPLOAD_INPUT_PARAMETER(src2) |
||||
UMAT_UPLOAD_INPUT_PARAMETER(weights1) |
||||
UMAT_UPLOAD_INPUT_PARAMETER(weights2) |
||||
UMAT_UPLOAD_OUTPUT_PARAMETER(dst) |
||||
} |
||||
|
||||
void Near(double eps = 0.0) |
||||
{ |
||||
EXPECT_MAT_NEAR(dst, udst, eps); |
||||
EXPECT_MAT_NEAR(dst_roi, udst_roi, eps); |
||||
} |
||||
}; |
||||
|
||||
OCL_TEST_P(BlendLinear, Accuracy) |
||||
{ |
||||
for (int i = 0; i < test_loop_times; ++i) |
||||
{ |
||||
random_roi(); |
||||
|
||||
OCL_OFF(cv::blendLinear(src1_roi, src2_roi, weights1_roi, weights2_roi, dst_roi)); |
||||
OCL_ON(cv::blendLinear(usrc1_roi, usrc2_roi, uweights1_roi, uweights2_roi, udst_roi)); |
||||
|
||||
Near(depth <= CV_32S ? 1.0 : 0.2); |
||||
} |
||||
} |
||||
|
||||
OCL_INSTANTIATE_TEST_CASE_P(ImgProc, BlendLinear, Combine(testing::Values(CV_8U, CV_32F), OCL_ALL_CHANNELS, Bool())); |
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif |
Loading…
Reference in new issue