2013-10-22 18:04:49 +08:00
|
|
|
/*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
|
|
|
|
//
|
2015-01-12 15:59:30 +08:00
|
|
|
// Copyright (C) 2014, Itseez Inc., all rights reserved.
|
2013-10-22 18:04:49 +08:00
|
|
|
// 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 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"
|
2014-08-01 22:11:20 +08:00
|
|
|
#include "opencl_kernels_core.hpp"
|
2018-01-16 22:33:06 +08:00
|
|
|
#include "umatrix.hpp"
|
2013-10-22 18:04:49 +08:00
|
|
|
|
2019-10-13 19:14:41 +08:00
|
|
|
#include <opencv2/core/utils/tls.hpp>
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
///////////////////////////////// UMat implementation ///////////////////////////////
|
|
|
|
|
|
|
|
namespace cv {
|
|
|
|
|
OpenCV-OpenCL interop (PR #4072):
Commits:
added new function, cv::ocl::attachContext(String& platformName, void* platformID, void* context, void* deviceID) which allow to attach externally created OpenCL context to OpenCV.
add definitions of clRetainDevice, clRetainContext funcs
removed definitions for clRetainContext, clRetainDevice
fixed build issue under Linux
fixed uninitialized vars, replace dbgassert in error handling
remove function which is not ready yet
add new function, cv::ocl::convertFromBuffer(int rows, int cols, int type, void* cl_mem_obj, UMat& dst, UMatUsageFlags usageFlags = cv::USAGE_DEFAULT) which attaches user allocated OpenCL clBuffer to UMat
uncommented clGetMemObjectInfo definition (otherwise prevent opencv build)
fixed build issue on linux and android
add step parameter to cv::ocl::convertFromBuffer func
suppress compile-time warning
added sample opencl-opencv interoperability (showcase for cv::ocl::convertFromBuffer func)
CMakeLists.txt modified to not create sample build script if OpenCL SDK not found in system
fixed build issue (apple opencl include dir and spaces in CMake file)
added call to clRetainContext for attachContext func and call to clRetainMemObject for convertFromBuffer func
uncommented clRetainMemObject definition
added comments and cleanup
add local path to cmake modules search dirs (instead of replacing)
remove REQUIRED for find_package call (sample build together with opencv). need to try standalone sample build
opencl-interop sample moved to standalone build
set minimum version requirement for sample's cmake to 3.1
put cmake_minimum_required under condition, so do not check if samples not builded
remove code dups for setSize, updateContinuityFlag, and finalizeHdr
commented out cmake_minimum_required(VERSION 3.1)
add safety check for cmake version
add convertFromImage func and update opencl-interop sample
uncommented clGetImageInfo defs
uncommented clEnqueueCopyImageToBuffer defs
fixed clEnqueueCopyImageToBuffer defs
add doxygen comments
remove doxygen @fn tag
try to restart buildbot
add doxygen comments to directx interop funcs
remove internal header, use fwd declarations in affected compile units instead
2015-05-28 04:22:33 +08:00
|
|
|
// forward decls, implementation is below in this file
|
|
|
|
void setSize(UMat& m, int _dims, const int* _sz, const size_t* _steps,
|
|
|
|
bool autoSteps = false);
|
|
|
|
|
|
|
|
void updateContinuityFlag(UMat& m);
|
|
|
|
void finalizeHdr(UMat& m);
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
UMatData::UMatData(const MatAllocator* allocator)
|
|
|
|
{
|
|
|
|
prevAllocator = currAllocator = allocator;
|
2015-09-03 22:18:59 +08:00
|
|
|
urefcount = refcount = mapcount = 0;
|
2013-10-22 18:04:49 +08:00
|
|
|
data = origdata = 0;
|
2015-01-02 08:33:40 +08:00
|
|
|
size = 0;
|
2018-09-21 23:12:35 +08:00
|
|
|
flags = static_cast<UMatData::MemoryFlag>(0);
|
2013-10-22 18:04:49 +08:00
|
|
|
handle = 0;
|
2013-10-23 03:34:16 +08:00
|
|
|
userdata = 0;
|
2014-02-10 20:34:45 +08:00
|
|
|
allocatorFlags_ = 0;
|
2015-09-08 09:06:04 +08:00
|
|
|
originalUMatData = NULL;
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
2013-12-01 07:12:19 +08:00
|
|
|
UMatData::~UMatData()
|
|
|
|
{
|
|
|
|
prevAllocator = currAllocator = 0;
|
|
|
|
urefcount = refcount = 0;
|
2015-09-03 22:18:59 +08:00
|
|
|
CV_Assert(mapcount == 0);
|
2013-12-01 07:12:19 +08:00
|
|
|
data = origdata = 0;
|
2015-01-02 08:33:40 +08:00
|
|
|
size = 0;
|
2020-08-21 03:05:50 +08:00
|
|
|
bool isAsyncCleanup = !!(flags & UMatData::ASYNC_CLEANUP);
|
2018-09-21 23:12:35 +08:00
|
|
|
flags = static_cast<UMatData::MemoryFlag>(0);
|
2013-12-01 07:12:19 +08:00
|
|
|
handle = 0;
|
|
|
|
userdata = 0;
|
2014-02-10 20:34:45 +08:00
|
|
|
allocatorFlags_ = 0;
|
2015-09-08 09:06:04 +08:00
|
|
|
if (originalUMatData)
|
|
|
|
{
|
|
|
|
bool showWarn = false;
|
2018-07-30 19:09:50 +08:00
|
|
|
UMatData* u = originalUMatData;
|
|
|
|
bool zero_Ref = CV_XADD(&(u->refcount), -1) == 1;
|
|
|
|
if (zero_Ref)
|
2015-09-08 09:06:04 +08:00
|
|
|
{
|
|
|
|
// simulate Mat::deallocate
|
|
|
|
if (u->mapcount != 0)
|
|
|
|
{
|
2015-12-01 04:45:48 +08:00
|
|
|
(u->currAllocator ? u->currAllocator : /* TODO allocator ? allocator :*/ Mat::getDefaultAllocator())->unmap(u);
|
2015-09-08 09:06:04 +08:00
|
|
|
}
|
|
|
|
else
|
|
|
|
{
|
|
|
|
// we don't do "map", so we can't do "unmap"
|
|
|
|
}
|
|
|
|
}
|
2018-07-30 19:09:50 +08:00
|
|
|
bool zero_URef = CV_XADD(&(u->urefcount), -1) == 1;
|
|
|
|
if (zero_Ref && !zero_URef)
|
|
|
|
showWarn = true;
|
|
|
|
if (zero_Ref && zero_URef) // oops, we need to free resources
|
2015-09-08 09:06:04 +08:00
|
|
|
{
|
2020-08-21 03:05:50 +08:00
|
|
|
showWarn = !isAsyncCleanup;
|
2015-09-08 09:06:04 +08:00
|
|
|
// simulate UMat::deallocate
|
|
|
|
u->currAllocator->deallocate(u);
|
|
|
|
}
|
|
|
|
#ifndef NDEBUG
|
|
|
|
if (showWarn)
|
|
|
|
{
|
|
|
|
static int warn_message_showed = 0;
|
|
|
|
if (warn_message_showed++ < 100)
|
|
|
|
{
|
|
|
|
fflush(stdout);
|
|
|
|
fprintf(stderr, "\n! OPENCV warning: getUMat()/getMat() call chain possible problem."
|
|
|
|
"\n! Base object is dead, while nested/derived object is still alive or processed."
|
|
|
|
"\n! Please check lifetime of UMat/Mat objects!\n");
|
|
|
|
fflush(stderr);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
#else
|
2018-09-07 19:33:52 +08:00
|
|
|
CV_UNUSED(showWarn);
|
2015-09-08 09:06:04 +08:00
|
|
|
#endif
|
|
|
|
originalUMatData = NULL;
|
|
|
|
}
|
2013-12-01 07:12:19 +08:00
|
|
|
}
|
|
|
|
|
2021-07-09 04:21:21 +08:00
|
|
|
#ifndef OPENCV_DISABLE_THREAD_SUPPORT
|
|
|
|
|
|
|
|
// it should be a prime number for the best hash function
|
|
|
|
enum { UMAT_NLOCKS = 31 };
|
|
|
|
static Mutex umatLocks[UMAT_NLOCKS];
|
|
|
|
|
2018-01-16 22:33:06 +08:00
|
|
|
static size_t getUMatDataLockIndex(const UMatData* u)
|
|
|
|
{
|
|
|
|
size_t idx = ((size_t)(void*)u) % UMAT_NLOCKS;
|
|
|
|
return idx;
|
|
|
|
}
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
void UMatData::lock()
|
|
|
|
{
|
2018-01-16 22:33:06 +08:00
|
|
|
size_t idx = getUMatDataLockIndex(this);
|
|
|
|
//printf("%d lock(%d)\n", cv::utils::getThreadID(), (int)idx);
|
|
|
|
umatLocks[idx].lock();
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
void UMatData::unlock()
|
|
|
|
{
|
2018-01-16 22:33:06 +08:00
|
|
|
size_t idx = getUMatDataLockIndex(this);
|
|
|
|
//printf("%d unlock(%d)\n", cv::utils::getThreadID(), (int)idx);
|
|
|
|
umatLocks[idx].unlock();
|
|
|
|
}
|
|
|
|
|
|
|
|
|
2018-02-09 18:58:29 +08:00
|
|
|
// Do not allow several lock() calls with different UMatData objects.
|
|
|
|
struct UMatDataAutoLocker
|
2018-01-16 22:33:06 +08:00
|
|
|
{
|
2018-02-09 18:58:29 +08:00
|
|
|
int usage_count;
|
|
|
|
UMatData* locked_objects[2];
|
|
|
|
UMatDataAutoLocker() : usage_count(0) { locked_objects[0] = NULL; locked_objects[1] = NULL; }
|
|
|
|
|
|
|
|
void lock(UMatData*& u1)
|
|
|
|
{
|
|
|
|
bool locked_1 = (u1 == locked_objects[0] || u1 == locked_objects[1]);
|
|
|
|
if (locked_1)
|
|
|
|
{
|
|
|
|
u1 = NULL;
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
CV_Assert(usage_count == 0); // UMatDataAutoLock can't be used multiple times from the same thread
|
|
|
|
usage_count = 1;
|
|
|
|
locked_objects[0] = u1;
|
|
|
|
u1->lock();
|
|
|
|
}
|
|
|
|
void lock(UMatData*& u1, UMatData*& u2)
|
|
|
|
{
|
|
|
|
bool locked_1 = (u1 == locked_objects[0] || u1 == locked_objects[1]);
|
|
|
|
bool locked_2 = (u2 == locked_objects[0] || u2 == locked_objects[1]);
|
|
|
|
if (locked_1)
|
|
|
|
u1 = NULL;
|
|
|
|
if (locked_2)
|
|
|
|
u2 = NULL;
|
|
|
|
if (locked_1 && locked_2)
|
|
|
|
return;
|
|
|
|
CV_Assert(usage_count == 0); // UMatDataAutoLock can't be used multiple times from the same thread
|
|
|
|
usage_count = 1;
|
|
|
|
locked_objects[0] = u1;
|
|
|
|
locked_objects[1] = u2;
|
|
|
|
if (u1)
|
|
|
|
u1->lock();
|
|
|
|
if (u2)
|
|
|
|
u2->lock();
|
|
|
|
}
|
|
|
|
void release(UMatData* u1, UMatData* u2)
|
|
|
|
{
|
|
|
|
if (u1 == NULL && u2 == NULL)
|
|
|
|
return;
|
|
|
|
CV_Assert(usage_count == 1);
|
|
|
|
usage_count = 0;
|
|
|
|
if (u1)
|
|
|
|
u1->unlock();
|
|
|
|
if (u2)
|
|
|
|
u2->unlock();
|
|
|
|
locked_objects[0] = NULL; locked_objects[1] = NULL;
|
|
|
|
}
|
2018-01-16 22:33:06 +08:00
|
|
|
};
|
2018-02-09 18:58:29 +08:00
|
|
|
static TLSData<UMatDataAutoLocker>& getUMatDataAutoLockerTLS()
|
2018-01-16 22:33:06 +08:00
|
|
|
{
|
2018-02-09 18:58:29 +08:00
|
|
|
CV_SINGLETON_LAZY_INIT_REF(TLSData<UMatDataAutoLocker>, new TLSData<UMatDataAutoLocker>());
|
2018-01-16 22:33:06 +08:00
|
|
|
}
|
2018-02-09 18:58:29 +08:00
|
|
|
static UMatDataAutoLocker& getUMatDataAutoLocker() { return getUMatDataAutoLockerTLS().getRef(); }
|
2018-01-16 22:33:06 +08:00
|
|
|
|
|
|
|
|
|
|
|
UMatDataAutoLock::UMatDataAutoLock(UMatData* u) : u1(u), u2(NULL)
|
|
|
|
{
|
2018-02-09 18:58:29 +08:00
|
|
|
getUMatDataAutoLocker().lock(u1);
|
2018-01-16 22:33:06 +08:00
|
|
|
}
|
|
|
|
UMatDataAutoLock::UMatDataAutoLock(UMatData* u1_, UMatData* u2_) : u1(u1_), u2(u2_)
|
|
|
|
{
|
|
|
|
if (getUMatDataLockIndex(u1) > getUMatDataLockIndex(u2))
|
|
|
|
{
|
|
|
|
std::swap(u1, u2);
|
|
|
|
}
|
2018-02-09 18:58:29 +08:00
|
|
|
getUMatDataAutoLocker().lock(u1, u2);
|
2018-01-16 22:33:06 +08:00
|
|
|
}
|
|
|
|
UMatDataAutoLock::~UMatDataAutoLock()
|
|
|
|
{
|
2018-02-09 18:58:29 +08:00
|
|
|
getUMatDataAutoLocker().release(u1, u2);
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
2021-07-09 04:21:21 +08:00
|
|
|
#else
|
|
|
|
|
|
|
|
void UMatData::lock()
|
|
|
|
{
|
|
|
|
// nothing in OPENCV_DISABLE_THREAD_SUPPORT mode
|
|
|
|
}
|
|
|
|
|
|
|
|
void UMatData::unlock()
|
|
|
|
{
|
|
|
|
// nothing in OPENCV_DISABLE_THREAD_SUPPORT mode
|
|
|
|
}
|
|
|
|
|
|
|
|
UMatDataAutoLock::UMatDataAutoLock(UMatData* u) : u1(u), u2(NULL)
|
|
|
|
{
|
|
|
|
// nothing in OPENCV_DISABLE_THREAD_SUPPORT mode
|
|
|
|
}
|
|
|
|
UMatDataAutoLock::UMatDataAutoLock(UMatData* u1_, UMatData* u2_) : u1(u1_), u2(u2_)
|
|
|
|
{
|
|
|
|
// nothing in OPENCV_DISABLE_THREAD_SUPPORT mode
|
|
|
|
}
|
|
|
|
UMatDataAutoLock::~UMatDataAutoLock()
|
|
|
|
{
|
|
|
|
// nothing in OPENCV_DISABLE_THREAD_SUPPORT mode
|
|
|
|
}
|
|
|
|
|
|
|
|
#endif // OPENCV_DISABLE_THREAD_SUPPORT
|
|
|
|
|
2020-10-22 06:47:56 +08:00
|
|
|
//////////////////////////////// UMat ////////////////////////////////
|
|
|
|
|
2021-03-02 05:42:13 +08:00
|
|
|
UMat::UMat(UMatUsageFlags _usageFlags) CV_NOEXCEPT
|
2020-10-22 06:47:56 +08:00
|
|
|
: flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(_usageFlags), u(0), offset(0), size(&rows)
|
|
|
|
{}
|
|
|
|
|
|
|
|
UMat::UMat(int _rows, int _cols, int _type, UMatUsageFlags _usageFlags)
|
|
|
|
: flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(_usageFlags), u(0), offset(0), size(&rows)
|
|
|
|
{
|
|
|
|
create(_rows, _cols, _type);
|
|
|
|
}
|
|
|
|
|
|
|
|
UMat::UMat(int _rows, int _cols, int _type, const Scalar& _s, UMatUsageFlags _usageFlags)
|
|
|
|
: flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(_usageFlags), u(0), offset(0), size(&rows)
|
|
|
|
{
|
|
|
|
create(_rows, _cols, _type);
|
|
|
|
*this = _s;
|
|
|
|
}
|
|
|
|
|
|
|
|
UMat::UMat(Size _sz, int _type, UMatUsageFlags _usageFlags)
|
|
|
|
: flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(_usageFlags), u(0), offset(0), size(&rows)
|
|
|
|
{
|
|
|
|
create( _sz.height, _sz.width, _type );
|
|
|
|
}
|
|
|
|
|
|
|
|
UMat::UMat(Size _sz, int _type, const Scalar& _s, UMatUsageFlags _usageFlags)
|
|
|
|
: flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(_usageFlags), u(0), offset(0), size(&rows)
|
|
|
|
{
|
|
|
|
create(_sz.height, _sz.width, _type);
|
|
|
|
*this = _s;
|
|
|
|
}
|
|
|
|
|
|
|
|
UMat::UMat(int _dims, const int* _sz, int _type, UMatUsageFlags _usageFlags)
|
|
|
|
: flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(_usageFlags), u(0), offset(0), size(&rows)
|
|
|
|
{
|
|
|
|
create(_dims, _sz, _type);
|
|
|
|
}
|
|
|
|
|
|
|
|
UMat::UMat(int _dims, const int* _sz, int _type, const Scalar& _s, UMatUsageFlags _usageFlags)
|
|
|
|
: flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(_usageFlags), u(0), offset(0), size(&rows)
|
|
|
|
{
|
|
|
|
create(_dims, _sz, _type);
|
|
|
|
*this = _s;
|
|
|
|
}
|
|
|
|
|
|
|
|
UMat::UMat(const UMat& m)
|
|
|
|
: flags(m.flags), dims(m.dims), rows(m.rows), cols(m.cols), allocator(m.allocator),
|
|
|
|
usageFlags(m.usageFlags), u(m.u), offset(m.offset), size(&rows)
|
|
|
|
{
|
|
|
|
addref();
|
|
|
|
if( m.dims <= 2 )
|
|
|
|
{
|
|
|
|
step[0] = m.step[0]; step[1] = m.step[1];
|
|
|
|
}
|
|
|
|
else
|
|
|
|
{
|
|
|
|
dims = 0;
|
|
|
|
copySize(m);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
UMat& UMat::operator=(const UMat& m)
|
|
|
|
{
|
|
|
|
if( this != &m )
|
|
|
|
{
|
|
|
|
const_cast<UMat&>(m).addref();
|
|
|
|
release();
|
|
|
|
flags = m.flags;
|
|
|
|
if( dims <= 2 && m.dims <= 2 )
|
|
|
|
{
|
|
|
|
dims = m.dims;
|
|
|
|
rows = m.rows;
|
|
|
|
cols = m.cols;
|
|
|
|
step[0] = m.step[0];
|
|
|
|
step[1] = m.step[1];
|
|
|
|
}
|
|
|
|
else
|
|
|
|
copySize(m);
|
|
|
|
allocator = m.allocator;
|
2021-05-05 04:29:15 +08:00
|
|
|
usageFlags = m.usageFlags;
|
2020-10-22 06:47:56 +08:00
|
|
|
u = m.u;
|
|
|
|
offset = m.offset;
|
|
|
|
}
|
|
|
|
return *this;
|
|
|
|
}
|
|
|
|
|
|
|
|
UMat UMat::clone() const
|
|
|
|
{
|
|
|
|
UMat m;
|
|
|
|
copyTo(m);
|
|
|
|
return m;
|
|
|
|
}
|
|
|
|
|
|
|
|
void UMat::assignTo(UMat& m, int _type) const
|
|
|
|
{
|
|
|
|
if( _type < 0 )
|
|
|
|
m = *this;
|
|
|
|
else
|
|
|
|
convertTo(m, _type);
|
|
|
|
}
|
|
|
|
|
|
|
|
void UMat::create(int _rows, int _cols, int _type, UMatUsageFlags _usageFlags)
|
|
|
|
{
|
|
|
|
int sz[] = {_rows, _cols};
|
|
|
|
create(2, sz, _type, _usageFlags);
|
|
|
|
}
|
|
|
|
|
|
|
|
void UMat::create(Size _sz, int _type, UMatUsageFlags _usageFlags)
|
|
|
|
{
|
|
|
|
create(_sz.height, _sz.width, _type, _usageFlags);
|
|
|
|
}
|
|
|
|
|
|
|
|
void UMat::addref()
|
|
|
|
{
|
|
|
|
if( u )
|
|
|
|
CV_XADD(&(u->urefcount), 1);
|
|
|
|
}
|
|
|
|
|
|
|
|
void UMat::release()
|
|
|
|
{
|
|
|
|
if( u && CV_XADD(&(u->urefcount), -1) == 1 )
|
|
|
|
deallocate();
|
|
|
|
for(int i = 0; i < dims; i++)
|
|
|
|
size.p[i] = 0;
|
|
|
|
u = 0;
|
|
|
|
}
|
|
|
|
|
|
|
|
bool UMat::empty() const
|
|
|
|
{
|
|
|
|
return u == 0 || total() == 0 || dims == 0;
|
|
|
|
}
|
|
|
|
|
|
|
|
size_t UMat::total() const
|
|
|
|
{
|
|
|
|
if( dims <= 2 )
|
|
|
|
return (size_t)rows * cols;
|
|
|
|
size_t p = 1;
|
|
|
|
for( int i = 0; i < dims; i++ )
|
|
|
|
p *= size[i];
|
|
|
|
return p;
|
|
|
|
}
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
|
2020-10-27 06:58:30 +08:00
|
|
|
UMat::UMat(UMat&& m)
|
|
|
|
: flags(m.flags), dims(m.dims), rows(m.rows), cols(m.cols), allocator(m.allocator),
|
|
|
|
usageFlags(m.usageFlags), u(m.u), offset(m.offset), size(&rows)
|
|
|
|
{
|
|
|
|
if (m.dims <= 2) // move new step/size info
|
|
|
|
{
|
|
|
|
step[0] = m.step[0];
|
|
|
|
step[1] = m.step[1];
|
|
|
|
}
|
|
|
|
else
|
|
|
|
{
|
|
|
|
CV_DbgAssert(m.step.p != m.step.buf);
|
|
|
|
step.p = m.step.p;
|
|
|
|
size.p = m.size.p;
|
|
|
|
m.step.p = m.step.buf;
|
|
|
|
m.size.p = &m.rows;
|
|
|
|
}
|
|
|
|
m.flags = MAGIC_VAL; m.dims = m.rows = m.cols = 0;
|
|
|
|
m.allocator = NULL;
|
|
|
|
m.u = NULL;
|
|
|
|
m.offset = 0;
|
|
|
|
}
|
|
|
|
|
|
|
|
UMat& UMat::operator=(UMat&& m)
|
|
|
|
{
|
|
|
|
if (this == &m)
|
|
|
|
return *this;
|
|
|
|
release();
|
|
|
|
flags = m.flags; dims = m.dims; rows = m.rows; cols = m.cols;
|
|
|
|
allocator = m.allocator; usageFlags = m.usageFlags;
|
|
|
|
u = m.u;
|
|
|
|
offset = m.offset;
|
|
|
|
if (step.p != step.buf) // release self step/size
|
|
|
|
{
|
|
|
|
fastFree(step.p);
|
|
|
|
step.p = step.buf;
|
|
|
|
size.p = &rows;
|
|
|
|
}
|
|
|
|
if (m.dims <= 2) // move new step/size info
|
|
|
|
{
|
|
|
|
step[0] = m.step[0];
|
|
|
|
step[1] = m.step[1];
|
|
|
|
}
|
|
|
|
else
|
|
|
|
{
|
|
|
|
CV_DbgAssert(m.step.p != m.step.buf);
|
|
|
|
step.p = m.step.p;
|
|
|
|
size.p = m.size.p;
|
|
|
|
m.step.p = m.step.buf;
|
|
|
|
m.size.p = &m.rows;
|
|
|
|
}
|
2021-05-05 04:29:15 +08:00
|
|
|
m.flags = MAGIC_VAL;
|
|
|
|
m.usageFlags = USAGE_DEFAULT;
|
|
|
|
m.dims = m.rows = m.cols = 0;
|
2020-10-27 06:58:30 +08:00
|
|
|
m.allocator = NULL;
|
|
|
|
m.u = NULL;
|
|
|
|
m.offset = 0;
|
|
|
|
return *this;
|
|
|
|
}
|
|
|
|
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
MatAllocator* UMat::getStdAllocator()
|
|
|
|
{
|
2014-04-02 02:20:13 +08:00
|
|
|
#ifdef HAVE_OPENCL
|
2017-11-24 22:34:02 +08:00
|
|
|
if (ocl::useOpenCL())
|
2013-12-01 07:12:19 +08:00
|
|
|
return ocl::getOpenCLAllocator();
|
2014-04-02 02:20:13 +08:00
|
|
|
#endif
|
2015-12-01 04:45:48 +08:00
|
|
|
return Mat::getDefaultAllocator();
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
void swap( UMat& a, UMat& b )
|
|
|
|
{
|
|
|
|
std::swap(a.flags, b.flags);
|
|
|
|
std::swap(a.dims, b.dims);
|
|
|
|
std::swap(a.rows, b.rows);
|
|
|
|
std::swap(a.cols, b.cols);
|
|
|
|
std::swap(a.allocator, b.allocator);
|
|
|
|
std::swap(a.u, b.u);
|
|
|
|
std::swap(a.offset, b.offset);
|
|
|
|
|
|
|
|
std::swap(a.size.p, b.size.p);
|
|
|
|
std::swap(a.step.p, b.step.p);
|
|
|
|
std::swap(a.step.buf[0], b.step.buf[0]);
|
|
|
|
std::swap(a.step.buf[1], b.step.buf[1]);
|
|
|
|
|
|
|
|
if( a.step.p == b.step.buf )
|
|
|
|
{
|
|
|
|
a.step.p = a.step.buf;
|
|
|
|
a.size.p = &a.rows;
|
|
|
|
}
|
|
|
|
|
|
|
|
if( b.step.p == a.step.buf )
|
|
|
|
{
|
|
|
|
b.step.p = b.step.buf;
|
|
|
|
b.size.p = &b.rows;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
|
OpenCV-OpenCL interop (PR #4072):
Commits:
added new function, cv::ocl::attachContext(String& platformName, void* platformID, void* context, void* deviceID) which allow to attach externally created OpenCL context to OpenCV.
add definitions of clRetainDevice, clRetainContext funcs
removed definitions for clRetainContext, clRetainDevice
fixed build issue under Linux
fixed uninitialized vars, replace dbgassert in error handling
remove function which is not ready yet
add new function, cv::ocl::convertFromBuffer(int rows, int cols, int type, void* cl_mem_obj, UMat& dst, UMatUsageFlags usageFlags = cv::USAGE_DEFAULT) which attaches user allocated OpenCL clBuffer to UMat
uncommented clGetMemObjectInfo definition (otherwise prevent opencv build)
fixed build issue on linux and android
add step parameter to cv::ocl::convertFromBuffer func
suppress compile-time warning
added sample opencl-opencv interoperability (showcase for cv::ocl::convertFromBuffer func)
CMakeLists.txt modified to not create sample build script if OpenCL SDK not found in system
fixed build issue (apple opencl include dir and spaces in CMake file)
added call to clRetainContext for attachContext func and call to clRetainMemObject for convertFromBuffer func
uncommented clRetainMemObject definition
added comments and cleanup
add local path to cmake modules search dirs (instead of replacing)
remove REQUIRED for find_package call (sample build together with opencv). need to try standalone sample build
opencl-interop sample moved to standalone build
set minimum version requirement for sample's cmake to 3.1
put cmake_minimum_required under condition, so do not check if samples not builded
remove code dups for setSize, updateContinuityFlag, and finalizeHdr
commented out cmake_minimum_required(VERSION 3.1)
add safety check for cmake version
add convertFromImage func and update opencl-interop sample
uncommented clGetImageInfo defs
uncommented clEnqueueCopyImageToBuffer defs
fixed clEnqueueCopyImageToBuffer defs
add doxygen comments
remove doxygen @fn tag
try to restart buildbot
add doxygen comments to directx interop funcs
remove internal header, use fwd declarations in affected compile units instead
2015-05-28 04:22:33 +08:00
|
|
|
void setSize( UMat& m, int _dims, const int* _sz,
|
|
|
|
const size_t* _steps, bool autoSteps )
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
CV_Assert( 0 <= _dims && _dims <= CV_MAX_DIM );
|
|
|
|
if( m.dims != _dims )
|
|
|
|
{
|
|
|
|
if( m.step.p != m.step.buf )
|
|
|
|
{
|
|
|
|
fastFree(m.step.p);
|
|
|
|
m.step.p = m.step.buf;
|
|
|
|
m.size.p = &m.rows;
|
|
|
|
}
|
|
|
|
if( _dims > 2 )
|
|
|
|
{
|
|
|
|
m.step.p = (size_t*)fastMalloc(_dims*sizeof(m.step.p[0]) + (_dims+1)*sizeof(m.size.p[0]));
|
|
|
|
m.size.p = (int*)(m.step.p + _dims) + 1;
|
|
|
|
m.size.p[-1] = _dims;
|
|
|
|
m.rows = m.cols = -1;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
m.dims = _dims;
|
|
|
|
if( !_sz )
|
|
|
|
return;
|
|
|
|
|
|
|
|
size_t esz = CV_ELEM_SIZE(m.flags), total = esz;
|
|
|
|
int i;
|
|
|
|
for( i = _dims-1; i >= 0; i-- )
|
|
|
|
{
|
|
|
|
int s = _sz[i];
|
|
|
|
CV_Assert( s >= 0 );
|
|
|
|
m.size.p[i] = s;
|
|
|
|
|
|
|
|
if( _steps )
|
|
|
|
m.step.p[i] = i < _dims-1 ? _steps[i] : esz;
|
|
|
|
else if( autoSteps )
|
|
|
|
{
|
|
|
|
m.step.p[i] = total;
|
|
|
|
int64 total1 = (int64)total*s;
|
|
|
|
if( (uint64)total1 != (size_t)total1 )
|
|
|
|
CV_Error( CV_StsOutOfRange, "The total matrix size does not fit to \"size_t\" type" );
|
|
|
|
total = (size_t)total1;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
if( _dims == 1 )
|
|
|
|
{
|
|
|
|
m.dims = 2;
|
|
|
|
m.cols = 1;
|
|
|
|
m.step[1] = esz;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
OpenCV-OpenCL interop (PR #4072):
Commits:
added new function, cv::ocl::attachContext(String& platformName, void* platformID, void* context, void* deviceID) which allow to attach externally created OpenCL context to OpenCV.
add definitions of clRetainDevice, clRetainContext funcs
removed definitions for clRetainContext, clRetainDevice
fixed build issue under Linux
fixed uninitialized vars, replace dbgassert in error handling
remove function which is not ready yet
add new function, cv::ocl::convertFromBuffer(int rows, int cols, int type, void* cl_mem_obj, UMat& dst, UMatUsageFlags usageFlags = cv::USAGE_DEFAULT) which attaches user allocated OpenCL clBuffer to UMat
uncommented clGetMemObjectInfo definition (otherwise prevent opencv build)
fixed build issue on linux and android
add step parameter to cv::ocl::convertFromBuffer func
suppress compile-time warning
added sample opencl-opencv interoperability (showcase for cv::ocl::convertFromBuffer func)
CMakeLists.txt modified to not create sample build script if OpenCL SDK not found in system
fixed build issue (apple opencl include dir and spaces in CMake file)
added call to clRetainContext for attachContext func and call to clRetainMemObject for convertFromBuffer func
uncommented clRetainMemObject definition
added comments and cleanup
add local path to cmake modules search dirs (instead of replacing)
remove REQUIRED for find_package call (sample build together with opencv). need to try standalone sample build
opencl-interop sample moved to standalone build
set minimum version requirement for sample's cmake to 3.1
put cmake_minimum_required under condition, so do not check if samples not builded
remove code dups for setSize, updateContinuityFlag, and finalizeHdr
commented out cmake_minimum_required(VERSION 3.1)
add safety check for cmake version
add convertFromImage func and update opencl-interop sample
uncommented clGetImageInfo defs
uncommented clEnqueueCopyImageToBuffer defs
fixed clEnqueueCopyImageToBuffer defs
add doxygen comments
remove doxygen @fn tag
try to restart buildbot
add doxygen comments to directx interop funcs
remove internal header, use fwd declarations in affected compile units instead
2015-05-28 04:22:33 +08:00
|
|
|
|
2018-05-14 20:29:14 +08:00
|
|
|
void UMat::updateContinuityFlag()
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
2018-05-14 20:29:14 +08:00
|
|
|
flags = cv::updateContinuityFlag(flags, dims, size.p, step.p);
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
|
OpenCV-OpenCL interop (PR #4072):
Commits:
added new function, cv::ocl::attachContext(String& platformName, void* platformID, void* context, void* deviceID) which allow to attach externally created OpenCL context to OpenCV.
add definitions of clRetainDevice, clRetainContext funcs
removed definitions for clRetainContext, clRetainDevice
fixed build issue under Linux
fixed uninitialized vars, replace dbgassert in error handling
remove function which is not ready yet
add new function, cv::ocl::convertFromBuffer(int rows, int cols, int type, void* cl_mem_obj, UMat& dst, UMatUsageFlags usageFlags = cv::USAGE_DEFAULT) which attaches user allocated OpenCL clBuffer to UMat
uncommented clGetMemObjectInfo definition (otherwise prevent opencv build)
fixed build issue on linux and android
add step parameter to cv::ocl::convertFromBuffer func
suppress compile-time warning
added sample opencl-opencv interoperability (showcase for cv::ocl::convertFromBuffer func)
CMakeLists.txt modified to not create sample build script if OpenCL SDK not found in system
fixed build issue (apple opencl include dir and spaces in CMake file)
added call to clRetainContext for attachContext func and call to clRetainMemObject for convertFromBuffer func
uncommented clRetainMemObject definition
added comments and cleanup
add local path to cmake modules search dirs (instead of replacing)
remove REQUIRED for find_package call (sample build together with opencv). need to try standalone sample build
opencl-interop sample moved to standalone build
set minimum version requirement for sample's cmake to 3.1
put cmake_minimum_required under condition, so do not check if samples not builded
remove code dups for setSize, updateContinuityFlag, and finalizeHdr
commented out cmake_minimum_required(VERSION 3.1)
add safety check for cmake version
add convertFromImage func and update opencl-interop sample
uncommented clGetImageInfo defs
uncommented clEnqueueCopyImageToBuffer defs
fixed clEnqueueCopyImageToBuffer defs
add doxygen comments
remove doxygen @fn tag
try to restart buildbot
add doxygen comments to directx interop funcs
remove internal header, use fwd declarations in affected compile units instead
2015-05-28 04:22:33 +08:00
|
|
|
void finalizeHdr(UMat& m)
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
2018-05-14 20:29:14 +08:00
|
|
|
m.updateContinuityFlag();
|
2013-10-22 18:04:49 +08:00
|
|
|
int d = m.dims;
|
|
|
|
if( d > 2 )
|
|
|
|
m.rows = m.cols = -1;
|
|
|
|
}
|
|
|
|
|
OpenCV-OpenCL interop (PR #4072):
Commits:
added new function, cv::ocl::attachContext(String& platformName, void* platformID, void* context, void* deviceID) which allow to attach externally created OpenCL context to OpenCV.
add definitions of clRetainDevice, clRetainContext funcs
removed definitions for clRetainContext, clRetainDevice
fixed build issue under Linux
fixed uninitialized vars, replace dbgassert in error handling
remove function which is not ready yet
add new function, cv::ocl::convertFromBuffer(int rows, int cols, int type, void* cl_mem_obj, UMat& dst, UMatUsageFlags usageFlags = cv::USAGE_DEFAULT) which attaches user allocated OpenCL clBuffer to UMat
uncommented clGetMemObjectInfo definition (otherwise prevent opencv build)
fixed build issue on linux and android
add step parameter to cv::ocl::convertFromBuffer func
suppress compile-time warning
added sample opencl-opencv interoperability (showcase for cv::ocl::convertFromBuffer func)
CMakeLists.txt modified to not create sample build script if OpenCL SDK not found in system
fixed build issue (apple opencl include dir and spaces in CMake file)
added call to clRetainContext for attachContext func and call to clRetainMemObject for convertFromBuffer func
uncommented clRetainMemObject definition
added comments and cleanup
add local path to cmake modules search dirs (instead of replacing)
remove REQUIRED for find_package call (sample build together with opencv). need to try standalone sample build
opencl-interop sample moved to standalone build
set minimum version requirement for sample's cmake to 3.1
put cmake_minimum_required under condition, so do not check if samples not builded
remove code dups for setSize, updateContinuityFlag, and finalizeHdr
commented out cmake_minimum_required(VERSION 3.1)
add safety check for cmake version
add convertFromImage func and update opencl-interop sample
uncommented clGetImageInfo defs
uncommented clEnqueueCopyImageToBuffer defs
fixed clEnqueueCopyImageToBuffer defs
add doxygen comments
remove doxygen @fn tag
try to restart buildbot
add doxygen comments to directx interop funcs
remove internal header, use fwd declarations in affected compile units instead
2015-05-28 04:22:33 +08:00
|
|
|
|
2018-09-21 23:12:35 +08:00
|
|
|
UMat Mat::getUMat(AccessFlag accessFlags, UMatUsageFlags usageFlags) const
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
UMat hdr;
|
2013-12-01 07:12:19 +08:00
|
|
|
if(!data)
|
2013-10-22 18:04:49 +08:00
|
|
|
return hdr;
|
Merge pull request #9114 from pengli:dnn_rebase
add libdnn acceleration to dnn module (#9114)
* import libdnn code
Signed-off-by: Li Peng <peng.li@intel.com>
* add convolution layer ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* add pooling layer ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* add softmax layer ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* add lrn layer ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* add innerproduct layer ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* add HAVE_OPENCL macro
Signed-off-by: Li Peng <peng.li@intel.com>
* fix for convolution ocl
Signed-off-by: Li Peng <peng.li@intel.com>
* enable getUMat() for multi-dimension Mat
Signed-off-by: Li Peng <peng.li@intel.com>
* use getUMat for ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* use CV_OCL_RUN macro
Signed-off-by: Li Peng <peng.li@intel.com>
* set OPENCL target when it is available
and disable fuseLayer for OCL target for the time being
Signed-off-by: Li Peng <peng.li@intel.com>
* fix innerproduct accuracy test
Signed-off-by: Li Peng <peng.li@intel.com>
* remove trailing space
Signed-off-by: Li Peng <peng.li@intel.com>
* Fixed tensorflow demo bug.
Root cause is that tensorflow has different algorithm with libdnn
to calculate convolution output dimension.
libdnn don't calculate output dimension anymore and just use one
passed in by config.
* split gemm ocl file
split it into gemm_buffer.cl and gemm_image.cl
Signed-off-by: Li Peng <peng.li@intel.com>
* Fix compile failure
Signed-off-by: Li Peng <peng.li@intel.com>
* check env flag for auto tuning
Signed-off-by: Li Peng <peng.li@intel.com>
* switch to new ocl kernels for softmax layer
Signed-off-by: Li Peng <peng.li@intel.com>
* update softmax layer
on some platform subgroup extension may not work well,
fallback to non subgroup ocl acceleration.
Signed-off-by: Li Peng <peng.li@intel.com>
* fallback to cpu path for fc layer with multi output
Signed-off-by: Li Peng <peng.li@intel.com>
* update output message
Signed-off-by: Li Peng <peng.li@intel.com>
* update fully connected layer
fallback to gemm API if libdnn return false
Signed-off-by: Li Peng <peng.li@intel.com>
* Add ReLU OCL implementation
* disable layer fusion for now
Signed-off-by: Li Peng <peng.li@intel.com>
* Add OCL implementation for concat layer
Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com>
* libdnn: update license and copyrights
Also refine libdnn coding style
Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com>
Signed-off-by: Li Peng <peng.li@intel.com>
* DNN: Don't link OpenCL library explicitly
* DNN: Make default preferableTarget to DNN_TARGET_CPU
User should set it to DNN_TARGET_OPENCL explicitly if want to
use OpenCL acceleration.
Also don't fusion when using DNN_TARGET_OPENCL
* DNN: refine coding style
* Add getOpenCLErrorString
* DNN: Use int32_t/uint32_t instread of alias
* Use namespace ocl4dnn to include libdnn things
* remove extra copyTo in softmax ocl path
Signed-off-by: Li Peng <peng.li@intel.com>
* update ReLU layer ocl path
Signed-off-by: Li Peng <peng.li@intel.com>
* Add prefer target property for layer class
It is used to indicate the target for layer forwarding,
either the default CPU target or OCL target.
Signed-off-by: Li Peng <peng.li@intel.com>
* Add cl_event based timer for cv::ocl
* Rename libdnn to ocl4dnn
Signed-off-by: Li Peng <peng.li@intel.com>
Signed-off-by: wzw <zhiwen.wu@intel.com>
* use UMat for ocl4dnn internal buffer
Remove allocateMemory which use clCreateBuffer directly
Signed-off-by: Li Peng <peng.li@intel.com>
Signed-off-by: wzw <zhiwen.wu@intel.com>
* enable buffer gemm in ocl4dnn innerproduct
Signed-off-by: Li Peng <peng.li@intel.com>
* replace int_tp globally for ocl4dnn kernels.
Signed-off-by: wzw <zhiwen.wu@intel.com>
Signed-off-by: Li Peng <peng.li@intel.com>
* create UMat for layer params
Signed-off-by: Li Peng <peng.li@intel.com>
* update sign ocl kernel
Signed-off-by: Li Peng <peng.li@intel.com>
* update image based gemm of inner product layer
Signed-off-by: Li Peng <peng.li@intel.com>
* remove buffer gemm of inner product layer
call cv::gemm API instead
Signed-off-by: Li Peng <peng.li@intel.com>
* change ocl4dnn forward parameter to UMat
Signed-off-by: Li Peng <peng.li@intel.com>
* Refine auto-tuning mechanism.
- Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory
for fine-tuned kernel configuration.
e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp,
the cache directory will be /home/tmp/spatialkernels/ on Linux.
- Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable
auto-tuning.
- OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling
for OpenCL command queue. This fix basic kernel get wrong running
time, i.e. 0ms.
- If creating cache directory failed, disable auto-tuning.
* Detect and create cache dir on windows
Signed-off-by: Li Peng <peng.li@intel.com>
* Refine gemm like convolution kernel.
Signed-off-by: Li Peng <peng.li@intel.com>
* Fix redundant swizzleWeights calling when use cached kernel config.
* Fix "out of resource" bug when auto-tuning too many kernels.
* replace cl_mem with UMat in ocl4dnnConvSpatial class
* OCL4DNN: reduce the tuning kernel candidate.
This patch could reduce 75% of the tuning candidates with less
than 2% performance impact for the final result.
Signed-off-by: Zhigang Gong <zhigang.gong@intel.com>
* replace cl_mem with umat in ocl4dnn convolution
Signed-off-by: Li Peng <peng.li@intel.com>
* remove weight_image_ of ocl4dnn inner product
Actually it is unused in the computation
Signed-off-by: Li Peng <peng.li@intel.com>
* Various fixes for ocl4dnn
1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel())
2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp
3. Code comments cleanup
4. ignore check on OCL cpu device
Signed-off-by: Li Peng <peng.li@intel.com>
* add build option for log softmax
Signed-off-by: Li Peng <peng.li@intel.com>
* remove unused ocl kernels in ocl4dnn
Signed-off-by: Li Peng <peng.li@intel.com>
* replace ocl4dnnSet with opencv setTo
Signed-off-by: Li Peng <peng.li@intel.com>
* replace ALIGN with cv::alignSize
Signed-off-by: Li Peng <peng.li@intel.com>
* check kernel build options
Signed-off-by: Li Peng <peng.li@intel.com>
* Handle program compilation fail properly.
* Use std::numeric_limits<float>::infinity() for large float number
* check ocl4dnn kernel compilation result
Signed-off-by: Li Peng <peng.li@intel.com>
* remove unused ctx_id
Signed-off-by: Li Peng <peng.li@intel.com>
* change clEnqueueNDRangeKernel to kernel.run()
Signed-off-by: Li Peng <peng.li@intel.com>
* change cl_mem to UMat in image based gemm
Signed-off-by: Li Peng <peng.li@intel.com>
* check intel subgroup support for lrn and pooling layer
Signed-off-by: Li Peng <peng.li@intel.com>
* Fix convolution bug if group is greater than 1
Signed-off-by: Li Peng <peng.li@intel.com>
* Set default layer preferableTarget to be DNN_TARGET_CPU
Signed-off-by: Li Peng <peng.li@intel.com>
* Add ocl perf test for convolution
Signed-off-by: Li Peng <peng.li@intel.com>
* Add more ocl accuracy test
Signed-off-by: Li Peng <peng.li@intel.com>
* replace cl_image with ocl::Image2D
Signed-off-by: Li Peng <peng.li@intel.com>
* Fix build failure in elementwise layer
Signed-off-by: Li Peng <peng.li@intel.com>
* use getUMat() to get blob data
Signed-off-by: Li Peng <peng.li@intel.com>
* replace cl_mem handle with ocl::KernelArg
Signed-off-by: Li Peng <peng.li@intel.com>
* dnn(build): don't use C++11, OPENCL_LIBRARIES fix
* dnn(ocl4dnn): remove unused OpenCL kernels
* dnn(ocl4dnn): extract OpenCL code into .cl files
* dnn(ocl4dnn): refine auto-tuning
Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING
environment variable to enable it.
Use a set of pre-tuned configs as default config if auto-tuning is disabled.
These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet,
AlexNet, ResNet-50
If default config is not suitable, use the first available kernel config
from the candidates. Candidate priority from high to low is gemm like kernel,
IDLF kernel, basick kernel.
* dnn(ocl4dnn): pooling doesn't use OpenCL subgroups
* dnn(ocl4dnn): fix perf test
OpenCV has default 3sec time limit for each performance test.
Warmup OpenCL backend outside of perf measurement loop.
* use ocl::KernelArg as much as possible
Signed-off-by: Li Peng <peng.li@intel.com>
* dnn(ocl4dnn): fix bias bug for gemm like kernel
* dnn(ocl4dnn): wrap cl_mem into UMat
Signed-off-by: Li Peng <peng.li@intel.com>
* dnn(ocl4dnn): Refine signature of kernel config
- Use more readable string as signture of kernel config
- Don't count device name and vendor in signature string
- Default kernel configurations are tuned for Intel GPU with
24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model.
* dnn(ocl4dnn): swap width/height in configuration
* dnn(ocl4dnn): enable configs for Intel OpenCL runtime only
* core: make configuration helper functions accessible from non-core modules
* dnn(ocl4dnn): update kernel auto-tuning behavior
Avoid unwanted creation of directories
* dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash
* dnn(ocl4dnn): remove redundant code
* dnn(ocl4dnn): Add more clear message for simd size dismatch.
* dnn(ocl4dnn): add const to const argument
Signed-off-by: Li Peng <peng.li@intel.com>
* dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel
* dnn(ocl4dnn): drop unused tuneLocalSize()
* dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method
* dnn(ocl4dnn): sanitize file names used for cache
* dnn(perf): enable Network tests with OpenCL
* dnn(ocl4dnn/conv): drop computeGlobalSize()
* dnn(ocl4dnn/conv): drop unused fields
* dnn(ocl4dnn/conv): simplify ctor
* dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL
* dnn(ocl4dnn/conv): drop unsupported double / untested half types
* dnn(ocl4dnn/conv): drop unused variable
* dnn(ocl4dnn/conv): alignSize/divUp
* dnn(ocl4dnn/conv): use enum values
* dnn(ocl4dnn): drop unused innerproduct variable
Signed-off-by: Li Peng <peng.li@intel.com>
* dnn(ocl4dnn): add an generic function to check cl option support
* dnn(ocl4dnn): run softmax subgroup version kernel first
Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 20:38:00 +08:00
|
|
|
if (data != datastart)
|
2015-09-08 09:06:04 +08:00
|
|
|
{
|
Merge pull request #9114 from pengli:dnn_rebase
add libdnn acceleration to dnn module (#9114)
* import libdnn code
Signed-off-by: Li Peng <peng.li@intel.com>
* add convolution layer ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* add pooling layer ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* add softmax layer ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* add lrn layer ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* add innerproduct layer ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* add HAVE_OPENCL macro
Signed-off-by: Li Peng <peng.li@intel.com>
* fix for convolution ocl
Signed-off-by: Li Peng <peng.li@intel.com>
* enable getUMat() for multi-dimension Mat
Signed-off-by: Li Peng <peng.li@intel.com>
* use getUMat for ocl acceleration
Signed-off-by: Li Peng <peng.li@intel.com>
* use CV_OCL_RUN macro
Signed-off-by: Li Peng <peng.li@intel.com>
* set OPENCL target when it is available
and disable fuseLayer for OCL target for the time being
Signed-off-by: Li Peng <peng.li@intel.com>
* fix innerproduct accuracy test
Signed-off-by: Li Peng <peng.li@intel.com>
* remove trailing space
Signed-off-by: Li Peng <peng.li@intel.com>
* Fixed tensorflow demo bug.
Root cause is that tensorflow has different algorithm with libdnn
to calculate convolution output dimension.
libdnn don't calculate output dimension anymore and just use one
passed in by config.
* split gemm ocl file
split it into gemm_buffer.cl and gemm_image.cl
Signed-off-by: Li Peng <peng.li@intel.com>
* Fix compile failure
Signed-off-by: Li Peng <peng.li@intel.com>
* check env flag for auto tuning
Signed-off-by: Li Peng <peng.li@intel.com>
* switch to new ocl kernels for softmax layer
Signed-off-by: Li Peng <peng.li@intel.com>
* update softmax layer
on some platform subgroup extension may not work well,
fallback to non subgroup ocl acceleration.
Signed-off-by: Li Peng <peng.li@intel.com>
* fallback to cpu path for fc layer with multi output
Signed-off-by: Li Peng <peng.li@intel.com>
* update output message
Signed-off-by: Li Peng <peng.li@intel.com>
* update fully connected layer
fallback to gemm API if libdnn return false
Signed-off-by: Li Peng <peng.li@intel.com>
* Add ReLU OCL implementation
* disable layer fusion for now
Signed-off-by: Li Peng <peng.li@intel.com>
* Add OCL implementation for concat layer
Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com>
* libdnn: update license and copyrights
Also refine libdnn coding style
Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com>
Signed-off-by: Li Peng <peng.li@intel.com>
* DNN: Don't link OpenCL library explicitly
* DNN: Make default preferableTarget to DNN_TARGET_CPU
User should set it to DNN_TARGET_OPENCL explicitly if want to
use OpenCL acceleration.
Also don't fusion when using DNN_TARGET_OPENCL
* DNN: refine coding style
* Add getOpenCLErrorString
* DNN: Use int32_t/uint32_t instread of alias
* Use namespace ocl4dnn to include libdnn things
* remove extra copyTo in softmax ocl path
Signed-off-by: Li Peng <peng.li@intel.com>
* update ReLU layer ocl path
Signed-off-by: Li Peng <peng.li@intel.com>
* Add prefer target property for layer class
It is used to indicate the target for layer forwarding,
either the default CPU target or OCL target.
Signed-off-by: Li Peng <peng.li@intel.com>
* Add cl_event based timer for cv::ocl
* Rename libdnn to ocl4dnn
Signed-off-by: Li Peng <peng.li@intel.com>
Signed-off-by: wzw <zhiwen.wu@intel.com>
* use UMat for ocl4dnn internal buffer
Remove allocateMemory which use clCreateBuffer directly
Signed-off-by: Li Peng <peng.li@intel.com>
Signed-off-by: wzw <zhiwen.wu@intel.com>
* enable buffer gemm in ocl4dnn innerproduct
Signed-off-by: Li Peng <peng.li@intel.com>
* replace int_tp globally for ocl4dnn kernels.
Signed-off-by: wzw <zhiwen.wu@intel.com>
Signed-off-by: Li Peng <peng.li@intel.com>
* create UMat for layer params
Signed-off-by: Li Peng <peng.li@intel.com>
* update sign ocl kernel
Signed-off-by: Li Peng <peng.li@intel.com>
* update image based gemm of inner product layer
Signed-off-by: Li Peng <peng.li@intel.com>
* remove buffer gemm of inner product layer
call cv::gemm API instead
Signed-off-by: Li Peng <peng.li@intel.com>
* change ocl4dnn forward parameter to UMat
Signed-off-by: Li Peng <peng.li@intel.com>
* Refine auto-tuning mechanism.
- Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory
for fine-tuned kernel configuration.
e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp,
the cache directory will be /home/tmp/spatialkernels/ on Linux.
- Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable
auto-tuning.
- OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling
for OpenCL command queue. This fix basic kernel get wrong running
time, i.e. 0ms.
- If creating cache directory failed, disable auto-tuning.
* Detect and create cache dir on windows
Signed-off-by: Li Peng <peng.li@intel.com>
* Refine gemm like convolution kernel.
Signed-off-by: Li Peng <peng.li@intel.com>
* Fix redundant swizzleWeights calling when use cached kernel config.
* Fix "out of resource" bug when auto-tuning too many kernels.
* replace cl_mem with UMat in ocl4dnnConvSpatial class
* OCL4DNN: reduce the tuning kernel candidate.
This patch could reduce 75% of the tuning candidates with less
than 2% performance impact for the final result.
Signed-off-by: Zhigang Gong <zhigang.gong@intel.com>
* replace cl_mem with umat in ocl4dnn convolution
Signed-off-by: Li Peng <peng.li@intel.com>
* remove weight_image_ of ocl4dnn inner product
Actually it is unused in the computation
Signed-off-by: Li Peng <peng.li@intel.com>
* Various fixes for ocl4dnn
1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel())
2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp
3. Code comments cleanup
4. ignore check on OCL cpu device
Signed-off-by: Li Peng <peng.li@intel.com>
* add build option for log softmax
Signed-off-by: Li Peng <peng.li@intel.com>
* remove unused ocl kernels in ocl4dnn
Signed-off-by: Li Peng <peng.li@intel.com>
* replace ocl4dnnSet with opencv setTo
Signed-off-by: Li Peng <peng.li@intel.com>
* replace ALIGN with cv::alignSize
Signed-off-by: Li Peng <peng.li@intel.com>
* check kernel build options
Signed-off-by: Li Peng <peng.li@intel.com>
* Handle program compilation fail properly.
* Use std::numeric_limits<float>::infinity() for large float number
* check ocl4dnn kernel compilation result
Signed-off-by: Li Peng <peng.li@intel.com>
* remove unused ctx_id
Signed-off-by: Li Peng <peng.li@intel.com>
* change clEnqueueNDRangeKernel to kernel.run()
Signed-off-by: Li Peng <peng.li@intel.com>
* change cl_mem to UMat in image based gemm
Signed-off-by: Li Peng <peng.li@intel.com>
* check intel subgroup support for lrn and pooling layer
Signed-off-by: Li Peng <peng.li@intel.com>
* Fix convolution bug if group is greater than 1
Signed-off-by: Li Peng <peng.li@intel.com>
* Set default layer preferableTarget to be DNN_TARGET_CPU
Signed-off-by: Li Peng <peng.li@intel.com>
* Add ocl perf test for convolution
Signed-off-by: Li Peng <peng.li@intel.com>
* Add more ocl accuracy test
Signed-off-by: Li Peng <peng.li@intel.com>
* replace cl_image with ocl::Image2D
Signed-off-by: Li Peng <peng.li@intel.com>
* Fix build failure in elementwise layer
Signed-off-by: Li Peng <peng.li@intel.com>
* use getUMat() to get blob data
Signed-off-by: Li Peng <peng.li@intel.com>
* replace cl_mem handle with ocl::KernelArg
Signed-off-by: Li Peng <peng.li@intel.com>
* dnn(build): don't use C++11, OPENCL_LIBRARIES fix
* dnn(ocl4dnn): remove unused OpenCL kernels
* dnn(ocl4dnn): extract OpenCL code into .cl files
* dnn(ocl4dnn): refine auto-tuning
Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING
environment variable to enable it.
Use a set of pre-tuned configs as default config if auto-tuning is disabled.
These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet,
AlexNet, ResNet-50
If default config is not suitable, use the first available kernel config
from the candidates. Candidate priority from high to low is gemm like kernel,
IDLF kernel, basick kernel.
* dnn(ocl4dnn): pooling doesn't use OpenCL subgroups
* dnn(ocl4dnn): fix perf test
OpenCV has default 3sec time limit for each performance test.
Warmup OpenCL backend outside of perf measurement loop.
* use ocl::KernelArg as much as possible
Signed-off-by: Li Peng <peng.li@intel.com>
* dnn(ocl4dnn): fix bias bug for gemm like kernel
* dnn(ocl4dnn): wrap cl_mem into UMat
Signed-off-by: Li Peng <peng.li@intel.com>
* dnn(ocl4dnn): Refine signature of kernel config
- Use more readable string as signture of kernel config
- Don't count device name and vendor in signature string
- Default kernel configurations are tuned for Intel GPU with
24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model.
* dnn(ocl4dnn): swap width/height in configuration
* dnn(ocl4dnn): enable configs for Intel OpenCL runtime only
* core: make configuration helper functions accessible from non-core modules
* dnn(ocl4dnn): update kernel auto-tuning behavior
Avoid unwanted creation of directories
* dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash
* dnn(ocl4dnn): remove redundant code
* dnn(ocl4dnn): Add more clear message for simd size dismatch.
* dnn(ocl4dnn): add const to const argument
Signed-off-by: Li Peng <peng.li@intel.com>
* dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel
* dnn(ocl4dnn): drop unused tuneLocalSize()
* dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method
* dnn(ocl4dnn): sanitize file names used for cache
* dnn(perf): enable Network tests with OpenCL
* dnn(ocl4dnn/conv): drop computeGlobalSize()
* dnn(ocl4dnn/conv): drop unused fields
* dnn(ocl4dnn/conv): simplify ctor
* dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL
* dnn(ocl4dnn/conv): drop unsupported double / untested half types
* dnn(ocl4dnn/conv): drop unused variable
* dnn(ocl4dnn/conv): alignSize/divUp
* dnn(ocl4dnn/conv): use enum values
* dnn(ocl4dnn): drop unused innerproduct variable
Signed-off-by: Li Peng <peng.li@intel.com>
* dnn(ocl4dnn): add an generic function to check cl option support
* dnn(ocl4dnn): run softmax subgroup version kernel first
Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 20:38:00 +08:00
|
|
|
Size wholeSize;
|
|
|
|
Point ofs;
|
|
|
|
locateROI(wholeSize, ofs);
|
|
|
|
Size sz(cols, rows);
|
|
|
|
if (ofs.x != 0 || ofs.y != 0)
|
|
|
|
{
|
|
|
|
Mat src = *this;
|
|
|
|
int dtop = ofs.y;
|
|
|
|
int dbottom = wholeSize.height - src.rows - ofs.y;
|
|
|
|
int dleft = ofs.x;
|
|
|
|
int dright = wholeSize.width - src.cols - ofs.x;
|
|
|
|
src.adjustROI(dtop, dbottom, dleft, dright);
|
|
|
|
return src.getUMat(accessFlags, usageFlags)(cv::Rect(ofs.x, ofs.y, sz.width, sz.height));
|
|
|
|
}
|
2015-09-08 09:06:04 +08:00
|
|
|
}
|
|
|
|
CV_Assert(data == datastart);
|
|
|
|
|
2015-07-25 00:10:31 +08:00
|
|
|
accessFlags |= ACCESS_RW;
|
2015-09-08 09:06:04 +08:00
|
|
|
UMatData* new_u = NULL;
|
2013-12-01 07:12:19 +08:00
|
|
|
{
|
2015-12-01 04:45:48 +08:00
|
|
|
MatAllocator *a = allocator, *a0 = getDefaultAllocator();
|
2013-12-01 07:12:19 +08:00
|
|
|
if(!a)
|
|
|
|
a = a0;
|
2015-09-08 09:06:04 +08:00
|
|
|
new_u = a->allocate(dims, size.p, type(), data, step.p, accessFlags, usageFlags);
|
2018-04-20 19:51:55 +08:00
|
|
|
new_u->originalUMatData = u;
|
2013-12-01 07:12:19 +08:00
|
|
|
}
|
2015-07-09 18:26:59 +08:00
|
|
|
bool allocated = false;
|
2018-11-09 00:46:25 +08:00
|
|
|
try
|
2015-07-09 18:26:59 +08:00
|
|
|
{
|
2015-09-08 09:06:04 +08:00
|
|
|
allocated = UMat::getStdAllocator()->allocate(new_u, accessFlags, usageFlags);
|
2015-07-09 18:26:59 +08:00
|
|
|
}
|
2018-11-09 00:46:25 +08:00
|
|
|
catch (const cv::Exception& e)
|
2015-07-09 18:26:59 +08:00
|
|
|
{
|
|
|
|
fprintf(stderr, "Exception: %s\n", e.what());
|
|
|
|
}
|
|
|
|
if (!allocated)
|
|
|
|
{
|
2015-12-01 04:45:48 +08:00
|
|
|
allocated = getDefaultAllocator()->allocate(new_u, accessFlags, usageFlags);
|
2015-07-09 18:26:59 +08:00
|
|
|
CV_Assert(allocated);
|
|
|
|
}
|
2015-09-08 09:06:04 +08:00
|
|
|
if (u != NULL)
|
|
|
|
{
|
|
|
|
#ifdef HAVE_OPENCL
|
|
|
|
if (ocl::useOpenCL() && new_u->currAllocator == ocl::getOpenCLAllocator())
|
|
|
|
{
|
|
|
|
CV_Assert(new_u->tempUMat());
|
|
|
|
}
|
|
|
|
#endif
|
|
|
|
CV_XADD(&(u->refcount), 1);
|
|
|
|
CV_XADD(&(u->urefcount), 1);
|
|
|
|
}
|
2021-08-25 00:56:25 +08:00
|
|
|
try
|
|
|
|
{
|
|
|
|
hdr.flags = flags;
|
2021-08-29 05:30:28 +08:00
|
|
|
hdr.usageFlags = usageFlags;
|
2021-08-25 00:56:25 +08:00
|
|
|
setSize(hdr, dims, size.p, step.p);
|
|
|
|
finalizeHdr(hdr);
|
|
|
|
hdr.u = new_u;
|
|
|
|
hdr.offset = 0; //data - datastart;
|
|
|
|
hdr.addref();
|
|
|
|
return hdr;
|
|
|
|
}
|
|
|
|
catch(...)
|
|
|
|
{
|
|
|
|
if (u != NULL)
|
|
|
|
{
|
|
|
|
CV_XADD(&(u->refcount), -1);
|
|
|
|
CV_XADD(&(u->urefcount), -1);
|
|
|
|
}
|
|
|
|
new_u->currAllocator->deallocate(new_u);
|
|
|
|
throw;
|
|
|
|
}
|
2021-08-29 05:30:28 +08:00
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
2014-02-10 20:34:45 +08:00
|
|
|
void UMat::create(int d, const int* _sizes, int _type, UMatUsageFlags _usageFlags)
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
int i;
|
|
|
|
CV_Assert(0 <= d && d <= CV_MAX_DIM && _sizes);
|
|
|
|
_type = CV_MAT_TYPE(_type);
|
|
|
|
|
2021-05-05 04:29:15 +08:00
|
|
|
// if param value is USAGE_DEFAULT by implicit default param value -or- explicit value
|
|
|
|
// ...then don't change the existing usageFlags
|
|
|
|
// it is not possible to change usage from non-default to USAGE_DEFAULT through create()
|
|
|
|
// ...instead must construct UMat()
|
|
|
|
if (_usageFlags == cv::USAGE_DEFAULT)
|
|
|
|
{
|
|
|
|
_usageFlags = usageFlags;
|
|
|
|
}
|
|
|
|
|
|
|
|
if( u && (d == dims || (d == 1 && dims <= 2)) && _type == type() && _usageFlags == usageFlags )
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
for( i = 0; i < d; i++ )
|
|
|
|
if( size[i] != _sizes[i] )
|
|
|
|
break;
|
|
|
|
if( i == d && (d > 1 || size[1] == 1))
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2016-01-26 18:59:10 +08:00
|
|
|
int _sizes_backup[CV_MAX_DIM]; // #5991
|
|
|
|
if (_sizes == (this->size.p))
|
|
|
|
{
|
|
|
|
for(i = 0; i < d; i++ )
|
|
|
|
_sizes_backup[i] = _sizes[i];
|
|
|
|
_sizes = _sizes_backup;
|
|
|
|
}
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
release();
|
2021-05-05 04:29:15 +08:00
|
|
|
usageFlags = _usageFlags;
|
2013-10-22 18:04:49 +08:00
|
|
|
if( d == 0 )
|
|
|
|
return;
|
|
|
|
flags = (_type & CV_MAT_TYPE_MASK) | MAGIC_VAL;
|
|
|
|
setSize(*this, d, _sizes, 0, true);
|
|
|
|
offset = 0;
|
|
|
|
|
|
|
|
if( total() > 0 )
|
|
|
|
{
|
|
|
|
MatAllocator *a = allocator, *a0 = getStdAllocator();
|
2015-07-09 18:26:59 +08:00
|
|
|
if (!a)
|
|
|
|
{
|
2013-10-22 18:04:49 +08:00
|
|
|
a = a0;
|
2015-12-01 04:45:48 +08:00
|
|
|
a0 = Mat::getDefaultAllocator();
|
2015-07-09 18:26:59 +08:00
|
|
|
}
|
2018-11-09 00:46:25 +08:00
|
|
|
try
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
2018-09-21 23:12:35 +08:00
|
|
|
u = a->allocate(dims, size, _type, 0, step.p, ACCESS_RW /* ignored */, usageFlags);
|
2013-10-22 18:04:49 +08:00
|
|
|
CV_Assert(u != 0);
|
|
|
|
}
|
2018-11-09 00:46:25 +08:00
|
|
|
catch(...)
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
if(a != a0)
|
2018-09-21 23:12:35 +08:00
|
|
|
u = a0->allocate(dims, size, _type, 0, step.p, ACCESS_RW /* ignored */, usageFlags);
|
2013-10-22 18:04:49 +08:00
|
|
|
CV_Assert(u != 0);
|
|
|
|
}
|
|
|
|
CV_Assert( step[dims-1] == (size_t)CV_ELEM_SIZE(flags) );
|
|
|
|
}
|
|
|
|
|
|
|
|
finalizeHdr(*this);
|
2013-12-16 20:46:36 +08:00
|
|
|
addref();
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
2016-12-15 02:56:43 +08:00
|
|
|
void UMat::create(const std::vector<int>& _sizes, int _type, UMatUsageFlags _usageFlags)
|
|
|
|
{
|
|
|
|
create((int)_sizes.size(), _sizes.data(), _type, _usageFlags);
|
|
|
|
}
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
void UMat::copySize(const UMat& m)
|
|
|
|
{
|
|
|
|
setSize(*this, m.dims, 0, 0);
|
|
|
|
for( int i = 0; i < dims; i++ )
|
|
|
|
{
|
|
|
|
size[i] = m.size[i];
|
|
|
|
step[i] = m.step[i];
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2013-12-01 07:12:19 +08:00
|
|
|
|
|
|
|
UMat::~UMat()
|
|
|
|
{
|
|
|
|
release();
|
|
|
|
if( step.p != step.buf )
|
|
|
|
fastFree(step.p);
|
|
|
|
}
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
void UMat::deallocate()
|
|
|
|
{
|
2017-07-25 17:26:53 +08:00
|
|
|
UMatData* u_ = u;
|
2013-12-16 20:46:36 +08:00
|
|
|
u = NULL;
|
2017-07-25 17:26:53 +08:00
|
|
|
u_->currAllocator->deallocate(u_);
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
|
|
|
|
UMat::UMat(const UMat& m, const Range& _rowRange, const Range& _colRange)
|
2014-02-21 18:04:01 +08:00
|
|
|
: flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(USAGE_DEFAULT), u(0), offset(0), size(&rows)
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
CV_Assert( m.dims >= 2 );
|
|
|
|
if( m.dims > 2 )
|
|
|
|
{
|
|
|
|
AutoBuffer<Range> rs(m.dims);
|
|
|
|
rs[0] = _rowRange;
|
|
|
|
rs[1] = _colRange;
|
|
|
|
for( int i = 2; i < m.dims; i++ )
|
|
|
|
rs[i] = Range::all();
|
2018-06-11 06:42:00 +08:00
|
|
|
*this = m(rs.data());
|
2013-10-22 18:04:49 +08:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
*this = m;
|
|
|
|
if( _rowRange != Range::all() && _rowRange != Range(0,rows) )
|
|
|
|
{
|
|
|
|
CV_Assert( 0 <= _rowRange.start && _rowRange.start <= _rowRange.end && _rowRange.end <= m.rows );
|
|
|
|
rows = _rowRange.size();
|
|
|
|
offset += step*_rowRange.start;
|
|
|
|
flags |= SUBMATRIX_FLAG;
|
|
|
|
}
|
|
|
|
|
|
|
|
if( _colRange != Range::all() && _colRange != Range(0,cols) )
|
|
|
|
{
|
|
|
|
CV_Assert( 0 <= _colRange.start && _colRange.start <= _colRange.end && _colRange.end <= m.cols );
|
|
|
|
cols = _colRange.size();
|
|
|
|
offset += _colRange.start*elemSize();
|
|
|
|
flags |= SUBMATRIX_FLAG;
|
|
|
|
}
|
|
|
|
|
2018-05-14 20:29:14 +08:00
|
|
|
updateContinuityFlag();
|
2013-10-22 18:04:49 +08:00
|
|
|
|
|
|
|
if( rows <= 0 || cols <= 0 )
|
|
|
|
{
|
|
|
|
release();
|
|
|
|
rows = cols = 0;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
|
|
UMat::UMat(const UMat& m, const Rect& roi)
|
|
|
|
: flags(m.flags), dims(2), rows(roi.height), cols(roi.width),
|
2014-02-21 18:04:01 +08:00
|
|
|
allocator(m.allocator), usageFlags(m.usageFlags), u(m.u), offset(m.offset + roi.y*m.step[0]), size(&rows)
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
CV_Assert( m.dims <= 2 );
|
|
|
|
|
|
|
|
size_t esz = CV_ELEM_SIZE(flags);
|
|
|
|
offset += roi.x*esz;
|
|
|
|
CV_Assert( 0 <= roi.x && 0 <= roi.width && roi.x + roi.width <= m.cols &&
|
|
|
|
0 <= roi.y && 0 <= roi.height && roi.y + roi.height <= m.rows );
|
|
|
|
if( roi.width < m.cols || roi.height < m.rows )
|
|
|
|
flags |= SUBMATRIX_FLAG;
|
|
|
|
|
|
|
|
step[0] = m.step[0]; step[1] = esz;
|
2018-05-14 20:29:14 +08:00
|
|
|
updateContinuityFlag();
|
2013-10-22 18:04:49 +08:00
|
|
|
|
2021-08-25 00:56:25 +08:00
|
|
|
addref();
|
2013-10-22 18:04:49 +08:00
|
|
|
if( rows <= 0 || cols <= 0 )
|
|
|
|
{
|
|
|
|
rows = cols = 0;
|
2021-08-25 00:56:25 +08:00
|
|
|
release();
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
|
|
UMat::UMat(const UMat& m, const Range* ranges)
|
2014-02-21 18:04:01 +08:00
|
|
|
: flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(USAGE_DEFAULT), u(0), offset(0), size(&rows)
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
int i, d = m.dims;
|
|
|
|
|
|
|
|
CV_Assert(ranges);
|
|
|
|
for( i = 0; i < d; i++ )
|
|
|
|
{
|
|
|
|
Range r = ranges[i];
|
|
|
|
CV_Assert( r == Range::all() || (0 <= r.start && r.start < r.end && r.end <= m.size[i]) );
|
|
|
|
}
|
|
|
|
*this = m;
|
|
|
|
for( i = 0; i < d; i++ )
|
|
|
|
{
|
|
|
|
Range r = ranges[i];
|
|
|
|
if( r != Range::all() && r != Range(0, size.p[i]))
|
|
|
|
{
|
|
|
|
size.p[i] = r.end - r.start;
|
|
|
|
offset += r.start*step.p[i];
|
|
|
|
flags |= SUBMATRIX_FLAG;
|
|
|
|
}
|
|
|
|
}
|
2018-05-14 20:29:14 +08:00
|
|
|
updateContinuityFlag();
|
2016-12-15 23:16:40 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
UMat::UMat(const UMat& m, const std::vector<Range>& ranges)
|
|
|
|
: flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(USAGE_DEFAULT), u(0), offset(0), size(&rows)
|
|
|
|
{
|
|
|
|
int i, d = m.dims;
|
|
|
|
|
|
|
|
CV_Assert((int)ranges.size() == d);
|
|
|
|
for (i = 0; i < d; i++)
|
|
|
|
{
|
|
|
|
Range r = ranges[i];
|
|
|
|
CV_Assert(r == Range::all() || (0 <= r.start && r.start < r.end && r.end <= m.size[i]));
|
|
|
|
}
|
|
|
|
*this = m;
|
|
|
|
for (i = 0; i < d; i++)
|
|
|
|
{
|
|
|
|
Range r = ranges[i];
|
|
|
|
if (r != Range::all() && r != Range(0, size.p[i]))
|
|
|
|
{
|
|
|
|
size.p[i] = r.end - r.start;
|
|
|
|
offset += r.start*step.p[i];
|
|
|
|
flags |= SUBMATRIX_FLAG;
|
|
|
|
}
|
|
|
|
}
|
2018-05-14 20:29:14 +08:00
|
|
|
updateContinuityFlag();
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
UMat UMat::diag(int d) const
|
|
|
|
{
|
|
|
|
CV_Assert( dims <= 2 );
|
|
|
|
UMat m = *this;
|
|
|
|
size_t esz = elemSize();
|
|
|
|
int len;
|
|
|
|
|
|
|
|
if( d >= 0 )
|
|
|
|
{
|
|
|
|
len = std::min(cols - d, rows);
|
|
|
|
m.offset += esz*d;
|
|
|
|
}
|
|
|
|
else
|
|
|
|
{
|
|
|
|
len = std::min(rows + d, cols);
|
|
|
|
m.offset -= step[0]*d;
|
|
|
|
}
|
|
|
|
CV_DbgAssert( len > 0 );
|
|
|
|
|
|
|
|
m.size[0] = m.rows = len;
|
|
|
|
m.size[1] = m.cols = 1;
|
|
|
|
m.step[0] += (len > 1 ? esz : 0);
|
|
|
|
|
2018-05-14 20:29:14 +08:00
|
|
|
m.updateContinuityFlag();
|
2013-10-22 18:04:49 +08:00
|
|
|
|
|
|
|
if( size() != Size(1,1) )
|
|
|
|
m.flags |= SUBMATRIX_FLAG;
|
|
|
|
|
|
|
|
return m;
|
|
|
|
}
|
|
|
|
|
|
|
|
void UMat::locateROI( Size& wholeSize, Point& ofs ) const
|
|
|
|
{
|
|
|
|
CV_Assert( dims <= 2 && step[0] > 0 );
|
|
|
|
size_t esz = elemSize(), minstep;
|
|
|
|
ptrdiff_t delta1 = (ptrdiff_t)offset, delta2 = (ptrdiff_t)u->size;
|
|
|
|
|
|
|
|
if( delta1 == 0 )
|
|
|
|
ofs.x = ofs.y = 0;
|
|
|
|
else
|
|
|
|
{
|
|
|
|
ofs.y = (int)(delta1/step[0]);
|
|
|
|
ofs.x = (int)((delta1 - step[0]*ofs.y)/esz);
|
|
|
|
CV_DbgAssert( offset == (size_t)(ofs.y*step[0] + ofs.x*esz) );
|
|
|
|
}
|
|
|
|
minstep = (ofs.x + cols)*esz;
|
|
|
|
wholeSize.height = (int)((delta2 - minstep)/step[0] + 1);
|
|
|
|
wholeSize.height = std::max(wholeSize.height, ofs.y + rows);
|
|
|
|
wholeSize.width = (int)((delta2 - step*(wholeSize.height-1))/esz);
|
|
|
|
wholeSize.width = std::max(wholeSize.width, ofs.x + cols);
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
|
|
UMat& UMat::adjustROI( int dtop, int dbottom, int dleft, int dright )
|
|
|
|
{
|
|
|
|
CV_Assert( dims <= 2 && step[0] > 0 );
|
|
|
|
Size wholeSize; Point ofs;
|
|
|
|
size_t esz = elemSize();
|
|
|
|
locateROI( wholeSize, ofs );
|
2017-02-22 19:00:22 +08:00
|
|
|
int row1 = std::min(std::max(ofs.y - dtop, 0), wholeSize.height), row2 = std::max(0, std::min(ofs.y + rows + dbottom, wholeSize.height));
|
|
|
|
int col1 = std::min(std::max(ofs.x - dleft, 0), wholeSize.width), col2 = std::max(0, std::min(ofs.x + cols + dright, wholeSize.width));
|
|
|
|
if(row1 > row2)
|
|
|
|
std::swap(row1, row2);
|
|
|
|
if(col1 > col2)
|
|
|
|
std::swap(col1, col2);
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
offset += (row1 - ofs.y)*step + (col1 - ofs.x)*esz;
|
|
|
|
rows = row2 - row1; cols = col2 - col1;
|
|
|
|
size.p[0] = rows; size.p[1] = cols;
|
2018-05-14 20:29:14 +08:00
|
|
|
updateContinuityFlag();
|
2013-10-22 18:04:49 +08:00
|
|
|
return *this;
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
|
|
UMat UMat::reshape(int new_cn, int new_rows) const
|
|
|
|
{
|
|
|
|
int cn = channels();
|
|
|
|
UMat hdr = *this;
|
|
|
|
|
|
|
|
if( dims > 2 && new_rows == 0 && new_cn != 0 && size[dims-1]*cn % new_cn == 0 )
|
|
|
|
{
|
|
|
|
hdr.flags = (hdr.flags & ~CV_MAT_CN_MASK) | ((new_cn-1) << CV_CN_SHIFT);
|
|
|
|
hdr.step[dims-1] = CV_ELEM_SIZE(hdr.flags);
|
|
|
|
hdr.size[dims-1] = hdr.size[dims-1]*cn / new_cn;
|
|
|
|
return hdr;
|
|
|
|
}
|
|
|
|
|
|
|
|
CV_Assert( dims <= 2 );
|
|
|
|
|
|
|
|
if( new_cn == 0 )
|
|
|
|
new_cn = cn;
|
|
|
|
|
|
|
|
int total_width = cols * cn;
|
|
|
|
|
|
|
|
if( (new_cn > total_width || total_width % new_cn != 0) && new_rows == 0 )
|
|
|
|
new_rows = rows * total_width / new_cn;
|
|
|
|
|
|
|
|
if( new_rows != 0 && new_rows != rows )
|
|
|
|
{
|
|
|
|
int total_size = total_width * rows;
|
|
|
|
if( !isContinuous() )
|
|
|
|
CV_Error( CV_BadStep,
|
|
|
|
"The matrix is not continuous, thus its number of rows can not be changed" );
|
|
|
|
|
|
|
|
if( (unsigned)new_rows > (unsigned)total_size )
|
|
|
|
CV_Error( CV_StsOutOfRange, "Bad new number of rows" );
|
|
|
|
|
|
|
|
total_width = total_size / new_rows;
|
|
|
|
|
|
|
|
if( total_width * new_rows != total_size )
|
|
|
|
CV_Error( CV_StsBadArg, "The total number of matrix elements "
|
|
|
|
"is not divisible by the new number of rows" );
|
|
|
|
|
|
|
|
hdr.rows = new_rows;
|
|
|
|
hdr.step[0] = total_width * elemSize1();
|
|
|
|
}
|
|
|
|
|
|
|
|
int new_width = total_width / new_cn;
|
|
|
|
|
|
|
|
if( new_width * new_cn != total_width )
|
|
|
|
CV_Error( CV_BadNumChannels,
|
|
|
|
"The total width is not divisible by the new number of channels" );
|
|
|
|
|
|
|
|
hdr.cols = new_width;
|
|
|
|
hdr.flags = (hdr.flags & ~CV_MAT_CN_MASK) | ((new_cn-1) << CV_CN_SHIFT);
|
|
|
|
hdr.step[1] = CV_ELEM_SIZE(hdr.flags);
|
|
|
|
return hdr;
|
|
|
|
}
|
|
|
|
|
2021-06-22 11:32:54 +08:00
|
|
|
UMat UMat::diag(const UMat& d, UMatUsageFlags usageFlags)
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
CV_Assert( d.cols == 1 || d.rows == 1 );
|
|
|
|
int len = d.rows + d.cols - 1;
|
2021-06-22 11:32:54 +08:00
|
|
|
UMat m(len, len, d.type(), Scalar(0), usageFlags);
|
2013-10-22 18:04:49 +08:00
|
|
|
UMat md = m.diag();
|
|
|
|
if( d.cols == 1 )
|
|
|
|
d.copyTo(md);
|
|
|
|
else
|
|
|
|
transpose(d, md);
|
|
|
|
return m;
|
|
|
|
}
|
|
|
|
|
|
|
|
int UMat::checkVector(int _elemChannels, int _depth, bool _requireContinuous) const
|
|
|
|
{
|
|
|
|
return (depth() == _depth || _depth <= 0) &&
|
|
|
|
(isContinuous() || !_requireContinuous) &&
|
|
|
|
((dims == 2 && (((rows == 1 || cols == 1) && channels() == _elemChannels) ||
|
|
|
|
(cols == _elemChannels && channels() == 1))) ||
|
|
|
|
(dims == 3 && channels() == 1 && size.p[2] == _elemChannels && (size.p[0] == 1 || size.p[1] == 1) &&
|
|
|
|
(isContinuous() || step.p[1] == step.p[2]*size.p[2])))
|
|
|
|
? (int)(total()*channels()/_elemChannels) : -1;
|
|
|
|
}
|
|
|
|
|
|
|
|
UMat UMat::reshape(int _cn, int _newndims, const int* _newsz) const
|
|
|
|
{
|
|
|
|
if(_newndims == dims)
|
|
|
|
{
|
|
|
|
if(_newsz == 0)
|
|
|
|
return reshape(_cn);
|
|
|
|
if(_newndims == 2)
|
|
|
|
return reshape(_cn, _newsz[0]);
|
|
|
|
}
|
|
|
|
|
2015-08-22 00:26:29 +08:00
|
|
|
if (isContinuous())
|
|
|
|
{
|
|
|
|
CV_Assert(_cn >= 0 && _newndims > 0 && _newndims <= CV_MAX_DIM && _newsz);
|
|
|
|
|
|
|
|
if (_cn == 0)
|
|
|
|
_cn = this->channels();
|
|
|
|
else
|
|
|
|
CV_Assert(_cn <= CV_CN_MAX);
|
|
|
|
|
|
|
|
size_t total_elem1_ref = this->total() * this->channels();
|
|
|
|
size_t total_elem1 = _cn;
|
|
|
|
|
|
|
|
AutoBuffer<int, 4> newsz_buf( (size_t)_newndims );
|
|
|
|
|
|
|
|
for (int i = 0; i < _newndims; i++)
|
|
|
|
{
|
|
|
|
CV_Assert(_newsz[i] >= 0);
|
|
|
|
|
|
|
|
if (_newsz[i] > 0)
|
|
|
|
newsz_buf[i] = _newsz[i];
|
|
|
|
else if (i < dims)
|
|
|
|
newsz_buf[i] = this->size[i];
|
|
|
|
else
|
|
|
|
CV_Error(CV_StsOutOfRange, "Copy dimension (which has zero size) is not present in source matrix");
|
|
|
|
|
|
|
|
total_elem1 *= (size_t)newsz_buf[i];
|
|
|
|
}
|
|
|
|
|
|
|
|
if (total_elem1 != total_elem1_ref)
|
|
|
|
CV_Error(CV_StsUnmatchedSizes, "Requested and source matrices have different count of elements");
|
|
|
|
|
|
|
|
UMat hdr = *this;
|
|
|
|
hdr.flags = (hdr.flags & ~CV_MAT_CN_MASK) | ((_cn-1) << CV_CN_SHIFT);
|
2018-06-11 06:42:00 +08:00
|
|
|
setSize(hdr, _newndims, newsz_buf.data(), NULL, true);
|
2015-08-22 00:26:29 +08:00
|
|
|
|
|
|
|
return hdr;
|
|
|
|
}
|
|
|
|
|
|
|
|
CV_Error(CV_StsNotImplemented, "Reshaping of n-dimensional non-continuous matrices is not supported yet");
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
2018-09-21 23:12:35 +08:00
|
|
|
Mat UMat::getMat(AccessFlag accessFlags) const
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
if(!u)
|
|
|
|
return Mat();
|
2015-07-25 00:10:31 +08:00
|
|
|
// TODO Support ACCESS_READ (ACCESS_WRITE) without unnecessary data transfers
|
|
|
|
accessFlags |= ACCESS_RW;
|
2015-09-03 22:18:59 +08:00
|
|
|
UMatDataAutoLock autolock(u);
|
2021-08-25 00:56:25 +08:00
|
|
|
try
|
2015-09-03 22:18:59 +08:00
|
|
|
{
|
2021-08-25 00:56:25 +08:00
|
|
|
if(CV_XADD(&u->refcount, 1) == 0)
|
|
|
|
u->currAllocator->map(u, accessFlags);
|
|
|
|
if (u->data != 0)
|
|
|
|
{
|
|
|
|
Mat hdr(dims, size.p, type(), u->data + offset, step.p);
|
|
|
|
hdr.flags = flags;
|
|
|
|
hdr.u = u;
|
|
|
|
hdr.datastart = u->data;
|
|
|
|
hdr.data = u->data + offset;
|
|
|
|
hdr.datalimit = hdr.dataend = u->data + u->size;
|
|
|
|
return hdr;
|
|
|
|
}
|
2015-09-03 22:18:59 +08:00
|
|
|
}
|
2021-08-25 00:56:25 +08:00
|
|
|
catch(...)
|
2015-09-03 22:18:59 +08:00
|
|
|
{
|
|
|
|
CV_XADD(&u->refcount, -1);
|
2021-08-25 00:56:25 +08:00
|
|
|
throw;
|
2015-09-03 22:18:59 +08:00
|
|
|
}
|
2021-08-25 00:56:25 +08:00
|
|
|
CV_XADD(&u->refcount, -1);
|
|
|
|
CV_Assert(u->data != 0 && "Error mapping of UMat to host memory.");
|
|
|
|
return Mat();
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
2018-09-21 23:12:35 +08:00
|
|
|
void* UMat::handle(AccessFlag accessFlags) const
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
if( !u )
|
|
|
|
return 0;
|
|
|
|
|
2015-09-08 09:06:04 +08:00
|
|
|
CV_Assert(u->refcount == 0);
|
|
|
|
CV_Assert(!u->deviceCopyObsolete() || u->copyOnMap());
|
|
|
|
if (u->deviceCopyObsolete())
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
u->currAllocator->unmap(u);
|
|
|
|
}
|
2014-08-07 16:31:50 +08:00
|
|
|
|
2018-09-21 23:12:35 +08:00
|
|
|
if (!!(accessFlags & ACCESS_WRITE))
|
2014-08-07 16:31:50 +08:00
|
|
|
u->markHostCopyObsolete(true);
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
return u->handle;
|
|
|
|
}
|
|
|
|
|
|
|
|
void UMat::ndoffset(size_t* ofs) const
|
|
|
|
{
|
|
|
|
// offset = step[0]*ofs[0] + step[1]*ofs[1] + step[2]*ofs[2] + ...;
|
2013-10-22 21:41:28 +08:00
|
|
|
size_t val = offset;
|
2013-10-22 18:04:49 +08:00
|
|
|
for( int i = 0; i < dims; i++ )
|
|
|
|
{
|
|
|
|
size_t s = step.p[i];
|
2013-10-22 21:41:28 +08:00
|
|
|
ofs[i] = val / s;
|
|
|
|
val -= ofs[i]*s;
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
void UMat::copyTo(OutputArray _dst) const
|
|
|
|
{
|
2018-09-14 05:35:26 +08:00
|
|
|
CV_INSTRUMENT_REGION();
|
2016-08-18 14:53:00 +08:00
|
|
|
|
2018-09-17 22:31:54 +08:00
|
|
|
#ifdef HAVE_CUDA
|
|
|
|
if (_dst.isGpuMat())
|
|
|
|
{
|
|
|
|
_dst.getGpuMat().upload(*this);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
#endif
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
int dtype = _dst.type();
|
|
|
|
if( _dst.fixedType() && dtype != type() )
|
|
|
|
{
|
|
|
|
CV_Assert( channels() == CV_MAT_CN(dtype) );
|
|
|
|
convertTo( _dst, dtype );
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
if( empty() )
|
|
|
|
{
|
|
|
|
_dst.release();
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2017-06-26 19:09:21 +08:00
|
|
|
size_t i, sz[CV_MAX_DIM] = {0}, srcofs[CV_MAX_DIM], dstofs[CV_MAX_DIM], esz = elemSize();
|
2013-10-22 18:04:49 +08:00
|
|
|
for( i = 0; i < (size_t)dims; i++ )
|
|
|
|
sz[i] = size.p[i];
|
2013-10-25 02:33:41 +08:00
|
|
|
sz[dims-1] *= esz;
|
2013-10-22 18:04:49 +08:00
|
|
|
ndoffset(srcofs);
|
2013-10-25 02:33:41 +08:00
|
|
|
srcofs[dims-1] *= esz;
|
2013-10-22 18:04:49 +08:00
|
|
|
|
2013-10-25 02:33:41 +08:00
|
|
|
_dst.create( dims, size.p, type() );
|
2014-02-02 19:33:10 +08:00
|
|
|
if( _dst.isUMat() )
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
|
|
|
UMat dst = _dst.getUMat();
|
2017-05-17 22:36:48 +08:00
|
|
|
CV_Assert(dst.u);
|
2013-12-17 18:14:04 +08:00
|
|
|
if( u == dst.u && dst.offset == offset )
|
2013-10-22 18:04:49 +08:00
|
|
|
return;
|
2014-02-02 19:33:10 +08:00
|
|
|
|
|
|
|
if (u->currAllocator == dst.u->currAllocator)
|
|
|
|
{
|
|
|
|
dst.ndoffset(dstofs);
|
|
|
|
dstofs[dims-1] *= esz;
|
|
|
|
u->currAllocator->copy(u, dst.u, dims, sz, srcofs, step.p, dstofs, dst.step.p, false);
|
|
|
|
return;
|
|
|
|
}
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
2014-02-02 19:33:10 +08:00
|
|
|
|
|
|
|
Mat dst = _dst.getMat();
|
2014-08-13 19:08:27 +08:00
|
|
|
u->currAllocator->download(u, dst.ptr(), dims, sz, srcofs, step.p, dst.step.p);
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
2013-12-21 00:10:43 +08:00
|
|
|
void UMat::copyTo(OutputArray _dst, InputArray _mask) const
|
|
|
|
{
|
2018-09-14 05:35:26 +08:00
|
|
|
CV_INSTRUMENT_REGION();
|
2016-08-18 14:53:00 +08:00
|
|
|
|
2013-12-21 00:10:43 +08:00
|
|
|
if( _mask.empty() )
|
|
|
|
{
|
|
|
|
copyTo(_dst);
|
|
|
|
return;
|
|
|
|
}
|
2014-04-02 02:20:13 +08:00
|
|
|
#ifdef HAVE_OPENCL
|
2013-12-21 00:10:43 +08:00
|
|
|
int cn = channels(), mtype = _mask.type(), mdepth = CV_MAT_DEPTH(mtype), mcn = CV_MAT_CN(mtype);
|
|
|
|
CV_Assert( mdepth == CV_8U && (mcn == 1 || mcn == cn) );
|
|
|
|
|
|
|
|
if (ocl::useOpenCL() && _dst.isUMat() && dims <= 2)
|
|
|
|
{
|
|
|
|
UMatData * prevu = _dst.getUMat().u;
|
|
|
|
_dst.create( dims, size, type() );
|
|
|
|
|
|
|
|
UMat dst = _dst.getUMat();
|
|
|
|
|
2014-05-30 22:01:49 +08:00
|
|
|
bool haveDstUninit = false;
|
2013-12-21 00:10:43 +08:00
|
|
|
if( prevu != dst.u ) // do not leave dst uninitialized
|
2014-05-30 22:01:49 +08:00
|
|
|
haveDstUninit = true;
|
2013-12-21 00:10:43 +08:00
|
|
|
|
2014-05-30 22:01:49 +08:00
|
|
|
String opts = format("-D COPY_TO_MASK -D T1=%s -D scn=%d -D mcn=%d%s",
|
|
|
|
ocl::memopTypeToStr(depth()), cn, mcn,
|
|
|
|
haveDstUninit ? " -D HAVE_DST_UNINIT" : "");
|
|
|
|
|
|
|
|
ocl::Kernel k("copyToMask", ocl::core::copyset_oclsrc, opts);
|
2013-12-21 00:10:43 +08:00
|
|
|
if (!k.empty())
|
|
|
|
{
|
2014-05-30 22:01:49 +08:00
|
|
|
k.args(ocl::KernelArg::ReadOnlyNoSize(*this),
|
|
|
|
ocl::KernelArg::ReadOnlyNoSize(_mask.getUMat()),
|
|
|
|
haveDstUninit ? ocl::KernelArg::WriteOnly(dst) :
|
|
|
|
ocl::KernelArg::ReadWrite(dst));
|
2013-12-21 00:10:43 +08:00
|
|
|
|
2015-10-16 22:10:00 +08:00
|
|
|
size_t globalsize[2] = { (size_t)cols, (size_t)rows };
|
2013-12-21 00:10:43 +08:00
|
|
|
if (k.run(2, globalsize, NULL, false))
|
2014-10-03 19:17:28 +08:00
|
|
|
{
|
|
|
|
CV_IMPL_ADD(CV_IMPL_OCL);
|
2013-12-21 00:10:43 +08:00
|
|
|
return;
|
2014-10-03 19:17:28 +08:00
|
|
|
}
|
2013-12-21 00:10:43 +08:00
|
|
|
}
|
|
|
|
}
|
2014-04-02 02:20:13 +08:00
|
|
|
#endif
|
2013-12-21 00:10:43 +08:00
|
|
|
Mat src = getMat(ACCESS_READ);
|
|
|
|
src.copyTo(_dst, _mask);
|
|
|
|
}
|
|
|
|
|
2016-08-18 14:53:00 +08:00
|
|
|
|
2024-01-26 17:56:52 +08:00
|
|
|
//
|
|
|
|
// void UMat::convertTo moved to convert.dispatch.cpp
|
|
|
|
//
|
2013-10-22 18:04:49 +08:00
|
|
|
|
2013-11-19 00:48:00 +08:00
|
|
|
UMat& UMat::setTo(InputArray _value, InputArray _mask)
|
|
|
|
{
|
2018-09-14 05:35:26 +08:00
|
|
|
CV_INSTRUMENT_REGION();
|
2016-08-18 14:53:00 +08:00
|
|
|
|
2013-11-19 00:48:00 +08:00
|
|
|
bool haveMask = !_mask.empty();
|
2014-04-02 02:20:13 +08:00
|
|
|
#ifdef HAVE_OPENCL
|
2015-01-05 22:56:46 +08:00
|
|
|
int tp = type(), cn = CV_MAT_CN(tp), d = CV_MAT_DEPTH(tp);
|
2014-04-02 02:20:13 +08:00
|
|
|
|
2014-02-12 23:29:18 +08:00
|
|
|
if( dims <= 2 && cn <= 4 && CV_MAT_DEPTH(tp) < CV_64F && ocl::useOpenCL() )
|
2013-11-19 00:48:00 +08:00
|
|
|
{
|
|
|
|
Mat value = _value.getMat();
|
|
|
|
CV_Assert( checkScalar(value, type(), _value.kind(), _InputArray::UMAT) );
|
2014-12-26 05:56:23 +08:00
|
|
|
int kercn = haveMask || cn == 3 ? cn : std::max(cn, ocl::predictOptimalVectorWidth(*this)),
|
2015-01-05 22:56:46 +08:00
|
|
|
kertp = CV_MAKE_TYPE(d, kercn);
|
2013-11-19 00:48:00 +08:00
|
|
|
|
2014-12-26 05:56:23 +08:00
|
|
|
double buf[16] = { 0, 0, 0, 0, 0, 0, 0, 0,
|
|
|
|
0, 0, 0, 0, 0, 0, 0, 0 };
|
|
|
|
convertAndUnrollScalar(value, tp, (uchar *)buf, kercn / cn);
|
|
|
|
|
|
|
|
int scalarcn = kercn == 3 ? 4 : kercn, rowsPerWI = ocl::Device::getDefault().isIntel() ? 4 : 1;
|
2014-05-27 16:07:26 +08:00
|
|
|
String opts = format("-D dstT=%s -D rowsPerWI=%d -D dstST=%s -D dstT1=%s -D cn=%d",
|
2014-12-26 05:56:23 +08:00
|
|
|
ocl::memopTypeToStr(kertp), rowsPerWI,
|
2015-01-05 22:56:46 +08:00
|
|
|
ocl::memopTypeToStr(CV_MAKETYPE(d, scalarcn)),
|
|
|
|
ocl::memopTypeToStr(d), kercn);
|
2013-11-19 00:48:00 +08:00
|
|
|
|
|
|
|
ocl::Kernel setK(haveMask ? "setMask" : "set", ocl::core::copyset_oclsrc, opts);
|
|
|
|
if( !setK.empty() )
|
|
|
|
{
|
2017-02-14 19:58:52 +08:00
|
|
|
ocl::KernelArg scalararg(ocl::KernelArg::CONSTANT, 0, 0, 0, buf, CV_ELEM_SIZE(d) * scalarcn);
|
2013-11-19 00:48:00 +08:00
|
|
|
UMat mask;
|
|
|
|
|
|
|
|
if( haveMask )
|
|
|
|
{
|
|
|
|
mask = _mask.getUMat();
|
2014-05-27 16:07:26 +08:00
|
|
|
CV_Assert( mask.size() == size() && mask.type() == CV_8UC1 );
|
|
|
|
ocl::KernelArg maskarg = ocl::KernelArg::ReadOnlyNoSize(mask),
|
|
|
|
dstarg = ocl::KernelArg::ReadWrite(*this);
|
2013-11-19 00:48:00 +08:00
|
|
|
setK.args(maskarg, dstarg, scalararg);
|
|
|
|
}
|
|
|
|
else
|
|
|
|
{
|
2014-12-26 05:56:23 +08:00
|
|
|
ocl::KernelArg dstarg = ocl::KernelArg::WriteOnly(*this, cn, kercn);
|
2013-11-19 00:48:00 +08:00
|
|
|
setK.args(dstarg, scalararg);
|
|
|
|
}
|
|
|
|
|
2015-10-16 22:10:00 +08:00
|
|
|
size_t globalsize[] = { (size_t)cols * cn / kercn, ((size_t)rows + rowsPerWI - 1) / rowsPerWI };
|
2014-05-27 16:07:26 +08:00
|
|
|
if( setK.run(2, globalsize, NULL, false) )
|
2014-10-03 19:17:28 +08:00
|
|
|
{
|
|
|
|
CV_IMPL_ADD(CV_IMPL_OCL);
|
2013-11-19 00:48:00 +08:00
|
|
|
return *this;
|
2014-10-03 19:17:28 +08:00
|
|
|
}
|
2013-11-19 00:48:00 +08:00
|
|
|
}
|
|
|
|
}
|
2014-04-02 02:20:13 +08:00
|
|
|
#endif
|
2013-11-19 00:48:00 +08:00
|
|
|
Mat m = getMat(haveMask ? ACCESS_RW : ACCESS_WRITE);
|
|
|
|
m.setTo(_value, _mask);
|
|
|
|
return *this;
|
|
|
|
}
|
|
|
|
|
2013-12-15 18:00:24 +08:00
|
|
|
UMat& UMat::operator = (const Scalar& s)
|
2013-10-22 18:04:49 +08:00
|
|
|
{
|
2013-12-15 18:00:24 +08:00
|
|
|
setTo(s);
|
2013-10-22 18:04:49 +08:00
|
|
|
return *this;
|
|
|
|
}
|
|
|
|
|
2014-02-05 23:10:02 +08:00
|
|
|
UMat UMat::t() const
|
|
|
|
{
|
|
|
|
UMat m;
|
|
|
|
transpose(*this, m);
|
|
|
|
return m;
|
|
|
|
}
|
|
|
|
|
2021-06-22 11:32:54 +08:00
|
|
|
UMat UMat::zeros(int rows, int cols, int type, UMatUsageFlags usageFlags)
|
2014-02-05 23:10:02 +08:00
|
|
|
{
|
2021-06-22 11:32:54 +08:00
|
|
|
return UMat(rows, cols, type, Scalar::all(0), usageFlags);
|
2014-02-05 23:10:02 +08:00
|
|
|
}
|
|
|
|
|
2021-06-22 11:32:54 +08:00
|
|
|
UMat UMat::zeros(Size size, int type, UMatUsageFlags usageFlags)
|
2014-02-05 23:10:02 +08:00
|
|
|
{
|
2021-06-22 11:32:54 +08:00
|
|
|
return UMat(size, type, Scalar::all(0), usageFlags);
|
2014-02-05 23:10:02 +08:00
|
|
|
}
|
|
|
|
|
2021-06-22 11:32:54 +08:00
|
|
|
UMat UMat::zeros(int ndims, const int* sz, int type, UMatUsageFlags usageFlags)
|
2014-02-05 23:10:02 +08:00
|
|
|
{
|
2021-06-22 11:32:54 +08:00
|
|
|
return UMat(ndims, sz, type, Scalar::all(0), usageFlags);
|
2014-02-05 23:10:02 +08:00
|
|
|
}
|
|
|
|
|
2021-06-22 11:32:54 +08:00
|
|
|
UMat UMat::ones(int rows, int cols, int type, UMatUsageFlags usageFlags)
|
2014-02-05 23:10:02 +08:00
|
|
|
{
|
2021-06-22 11:32:54 +08:00
|
|
|
return UMat(rows, cols, type, Scalar(1), usageFlags);
|
2014-02-05 23:10:02 +08:00
|
|
|
}
|
|
|
|
|
2021-06-22 11:32:54 +08:00
|
|
|
UMat UMat::ones(Size size, int type, UMatUsageFlags usageFlags)
|
2014-02-05 23:10:02 +08:00
|
|
|
{
|
2021-06-22 11:32:54 +08:00
|
|
|
return UMat(size, type, Scalar(1), usageFlags);
|
2014-02-05 23:10:02 +08:00
|
|
|
}
|
|
|
|
|
2021-06-22 11:32:54 +08:00
|
|
|
UMat UMat::ones(int ndims, const int* sz, int type, UMatUsageFlags usageFlags)
|
2014-02-05 23:10:02 +08:00
|
|
|
{
|
2021-06-22 11:32:54 +08:00
|
|
|
return UMat(ndims, sz, type, Scalar(1), usageFlags);
|
2014-02-05 23:10:02 +08:00
|
|
|
}
|
|
|
|
|
2013-10-22 18:04:49 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
/* End of file. */
|