diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index 919452dc98..689b57d885 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -41,6 +41,7 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels.hpp" /****************************************************************************************\ * [scaled] Identity matrix initialization * @@ -2368,10 +2369,35 @@ void cv::vconcat(InputArray _src, OutputArray dst) } //////////////////////////////////////// set identity //////////////////////////////////////////// + +namespace cv { + +static bool ocl_setIdentity( InputOutputArray _m, const Scalar& s ) +{ + int type = _m.type(), cn = CV_MAT_CN(type); + if (cn == 3) + return false; + + UMat m = _m.getUMat(); + + ocl::Kernel k("setIdentity", ocl::core::set_identity_oclsrc, + format("-D T=%s", ocl::memopTypeToStr(type))); + k.args(ocl::KernelArg::WriteOnly(m), ocl::KernelArg::Constant(Mat(1, 1, type, s))); + + size_t globalsize[2] = { m.cols, m.rows }; + return k.run(2, globalsize, NULL, false); +} + +} + void cv::setIdentity( InputOutputArray _m, const Scalar& s ) { + CV_Assert( _m.dims() <= 2 ); + + if (ocl::useOpenCL() && _m.isUMat() && ocl_setIdentity(_m, s)) + return; + Mat m = _m.getMat(); - CV_Assert( m.dims <= 2 ); int i, j, rows = m.rows, cols = m.cols, type = m.type(); if( type == CV_32FC1 ) diff --git a/modules/core/src/opencl/set_identity.cl b/modules/core/src/opencl/set_identity.cl new file mode 100644 index 0000000000..de8caaf85b --- /dev/null +++ b/modules/core/src/opencl/set_identity.cl @@ -0,0 +1,59 @@ +/*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 +// Jin Ma jin@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*/ + +__kernel void setIdentity(__global uchar * srcptr, int src_step, int src_offset, int rows, int cols, + T scalar) +{ + 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(T)); + __global T * src = (__global T *)(srcptr + src_index); + + src[0] = x == y ? scalar : (T)(0); + } +}