diff --git a/modules/core/src/opencl/convert.cl b/modules/core/src/opencl/convert.cl new file mode 100644 index 0000000000..6c2d16cff1 --- /dev/null +++ b/modules/core/src/opencl/convert.cl @@ -0,0 +1,71 @@ +/*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, Institute Of Software Chinese Academy Of Science, all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2013, OpenCV Foundation, all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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 copyright holders 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 convertTo(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, + float alpha, float beta ) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (x < dst_cols && y < dst_rows) + { + int src_index = mad24(y, src_step, src_offset + x * (int)sizeof(srcT) ); + int dst_index = mad24(y, dst_step, dst_offset + x * (int)sizeof(dstT) ); + + __global const srcT * src = (__global const srcT *)(srcptr + src_index); + __global dstT * dst = (__global dstT *)(dstptr + dst_index); + + dst[0] = convertToDT( src[0] * alpha + beta ); + } +} diff --git a/modules/core/src/umatrix.cpp b/modules/core/src/umatrix.cpp index 2d97a04b4a..12483877bc 100644 --- a/modules/core/src/umatrix.cpp +++ b/modules/core/src/umatrix.cpp @@ -658,16 +658,56 @@ void UMat::copyTo(OutputArray _dst) const } } -void UMat::convertTo(OutputArray, int, double, double) const +void UMat::convertTo(OutputArray _dst, int _type, double alpha, double beta) const { - CV_Error(Error::StsNotImplemented, ""); + bool noScale = std::fabs(alpha - 1) < DBL_EPSILON && std::fabs(beta) < DBL_EPSILON; + int stype = type(), cn = CV_MAT_CN(stype); + + if( _type < 0 ) + _type = _dst.fixedType() ? _dst.type() : stype; + else + _type = CV_MAKETYPE(CV_MAT_DEPTH(_type), cn); + + int sdepth = CV_MAT_DEPTH(stype), ddepth = CV_MAT_DEPTH(_type); + if( sdepth == ddepth && noScale ) + { + copyTo(_dst); + return; + } + + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + bool needDouble = sdepth == CV_64F || ddepth == CV_64F; + if( dims <= 2 && cn && _dst.isUMat() && ocl::useOpenCL() && + ((needDouble && doubleSupport) || !needDouble) ) + { + char cvt[40]; + ocl::Kernel k("convertTo", ocl::core::convert_oclsrc, + format("-D srcT=%s -D dstT=%s -D convertToDT=%s%s", ocl::typeToStr(sdepth), + ocl::typeToStr(ddepth), ocl::convertTypeStr(CV_32F, ddepth, 1, cvt), + doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + if (!k.empty()) + { + _dst.create( size(), _type ); + UMat dst = _dst.getUMat(); + + float alphaf = (float)alpha, betaf = (float)beta; + k.args(ocl::KernelArg::ReadOnlyNoSize(*this), ocl::KernelArg::WriteOnly(dst, cn), alphaf, betaf); + + size_t globalsize[2] = { dst.cols * cn, dst.rows }; + if (k.run(2, globalsize, NULL, false)) + return; + } + } + + Mat m = getMat(ACCESS_READ); + m.convertTo(_dst, _type, alpha, beta); } UMat& UMat::setTo(InputArray _value, InputArray _mask) { bool haveMask = !_mask.empty(); int tp = type(), cn = CV_MAT_CN(tp); - if( dims <= 2 && cn <= 4 && ocl::useOpenCL() ) + if( dims <= 2 && cn <= 4 && cn != 3 && ocl::useOpenCL() ) { Mat value = _value.getMat(); CV_Assert( checkScalar(value, type(), _value.kind(), _InputArray::UMAT) ); @@ -707,9 +747,9 @@ UMat& UMat::setTo(InputArray _value, InputArray _mask) return *this; } -UMat& UMat::operator = (const Scalar&) +UMat& UMat::operator = (const Scalar& s) { - CV_Error(Error::StsNotImplemented, ""); + setTo(s); return *this; } diff --git a/modules/core/test/ocl/test_matrix_operation.cpp b/modules/core/test/ocl/test_matrix_operation.cpp new file mode 100644 index 0000000000..599eb065ad --- /dev/null +++ b/modules/core/test/ocl/test_matrix_operation.cpp @@ -0,0 +1,111 @@ +/*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, Institute Of Software Chinese Academy Of Science, all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Jia Haipeng, jiahaipeng95@gmail.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 "opencv2/ts/ocl_test.hpp" + +#ifdef HAVE_OPENCL + +namespace cvtest { +namespace ocl { + +////////////////////////////////converto///////////////////////////////////////////////// + +PARAM_TEST_CASE(MatrixTestBase, MatDepth, MatDepth, Channels, bool) +{ + int src_depth, cn, dstType; + bool use_roi; + + TEST_DECLARE_INPUT_PARAMETER(src) + TEST_DECLARE_OUTPUT_PARAMETER(dst) + + virtual void SetUp() + { + src_depth = GET_PARAM(0); + cn = GET_PARAM(2); + dstType = CV_MAKE_TYPE(GET_PARAM(1), cn); + + use_roi = GET_PARAM(3); + } + + virtual void generateTestData() + { + Size roiSize = randomSize(1, MAX_VALUE); + Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src, src_roi, roiSize, srcBorder, CV_MAKE_TYPE(src_depth, cn), -MAX_VALUE, MAX_VALUE); + + 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) + } +}; + +typedef MatrixTestBase ConvertTo; + +OCL_TEST_P(ConvertTo, Accuracy) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + double alpha = randomDouble(-4, 4), beta = randomDouble(-4, 4); + + OCL_OFF(src_roi.convertTo(dst_roi, dstType, alpha, beta)); + OCL_ON(usrc_roi.convertTo(udst_roi, dstType, alpha, beta)); + + double eps = src_depth >= CV_32F || CV_MAT_DEPTH(dstType) >= CV_32F ? 1e-4 : 1; + OCL_EXPECT_MATS_NEAR(dst, eps); + } +} + +OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, ConvertTo, Combine( + OCL_ALL_DEPTHS, OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); + +} } // namespace cvtest::ocl + +#endif diff --git a/modules/core/test/test_umat.cpp b/modules/core/test/test_umat.cpp index 653dc49fb9..1dfe1d79a9 100644 --- a/modules/core/test/test_umat.cpp +++ b/modules/core/test/test_umat.cpp @@ -43,16 +43,13 @@ #include #include -#include -#include -#include -#include #include "opencv2/core/ocl.hpp" using namespace cv; using namespace std; -class CV_UMatTest : public cvtest::BaseTest +class CV_UMatTest : + public cvtest::BaseTest { public: CV_UMatTest() {} @@ -204,37 +201,37 @@ TEST(Core_UMat, base) { CV_UMatTest test; test.safe_run(); } TEST(Core_UMat, getUMat) { { - int a[3] = { 1, 2, 3 }; - Mat m = Mat(1, 1, CV_32SC3, a); - UMat u = m.getUMat(ACCESS_READ); - EXPECT_NE((void*)NULL, u.u); + int a[3] = { 1, 2, 3 }; + Mat m = Mat(1, 1, CV_32SC3, a); + UMat u = m.getUMat(ACCESS_READ); + EXPECT_NE((void*)NULL, u.u); } { - Mat m(10, 10, CV_8UC1), ref; - for (int y = 0; y < m.rows; ++y) - { - uchar * const ptr = m.ptr(y); - for (int x = 0; x < m.cols; ++x) - ptr[x] = (uchar)(x + y * 2); - } + Mat m(10, 10, CV_8UC1), ref; + for (int y = 0; y < m.rows; ++y) + { + uchar * const ptr = m.ptr(y); + for (int x = 0; x < m.cols; ++x) + ptr[x] = (uchar)(x + y * 2); + } - ref = m.clone(); - Rect r(1, 1, 8, 8); - ref(r).setTo(17); + ref = m.clone(); + Rect r(1, 1, 8, 8); + ref(r).setTo(17); - { - UMat um = m(r).getUMat(ACCESS_WRITE); - um.setTo(17); - } + { + UMat um = m(r).getUMat(ACCESS_WRITE); + um.setTo(17); + } - double err = norm(m, ref, NORM_INF); - if(err > 0) - { - std::cout << "m: " << m << std::endl; - std::cout << "ref: " << ref << std::endl; - } - EXPECT_EQ(err, 0.); + double err = norm(m, ref, NORM_INF); + if (err > 0) + { + std::cout << "m: " << std::endl << m << std::endl; + std::cout << "ref: " << std::endl << ref << std::endl; + } + EXPECT_EQ(0., err); } }