mirror of
https://github.com/opencv/opencv.git
synced 2026-07-30 07:43:03 +04:00
Merge remote-tracking branch 'origin/2.4'
Original pull requests: #996 from jet47:gpu-nvcuvid-libraries #995 from jet47:fix-bug-2985 #999 from snosov1:unreliable-results-fix #1005 from alekcac:doc_fix #1004 from jet47:fix-bug-3068 #987 from jet47:bug-3085-fix #969 from pengx17:2.4_binary_cache #929 from dominikrose:mingw-libdc1394-2-windows #1000 from ivan-korolev:fix_sift_bug_2892 #1001 from ivan-korolev:fix_stitching_bug_2405 #998 from asmorkalov:android_cmake_mips_fix #993 from ivan-korolev:fix_videostab_bug_3023 #988 from snosov1:3071-fix #986 from pengx17:2.4_initiated_context #982 from pengx17:2.4_fix_two_bugs #981 from SeninAndrew:ximea_camera_support_fix #991 from asmorkalov:android_javadoc_fix #972 from jet47:mog2-params-bug-2168 #980 from SpecLad:include-config #973 from pengx17:2.4_oclclahe #903 from aks2:2.4 #968 from asmorkalov:android_na_cproj_fix #971 from SpecLad:matchers-ctor #970 from asmorkalov:dshow_valid_check_fix #965 from apavlenko:fix_java_empty_mats Conflicts: cmake/OpenCVModule.cmake modules/core/src/matmul.cpp modules/gpu/CMakeLists.txt modules/ocl/include/opencv2/ocl/ocl.hpp modules/ocl/perf/perf_imgproc.cpp modules/ocl/src/imgproc.cpp modules/ocl/src/initialization.cpp modules/stitching/src/matchers.cpp modules/video/src/video_init.cpp modules/videostab/src/global_motion.cpp
This commit is contained in:
@@ -1741,7 +1741,7 @@ Returns the depth of a matrix element.
|
||||
|
||||
.. ocv:function:: int Mat::depth() const
|
||||
|
||||
The method returns the identifier of the matrix element depth (the type of each individual channel). For example, for a 16-bit signed 3-channel array, the method returns ``CV_16S`` . A complete list of matrix types contains the following values:
|
||||
The method returns the identifier of the matrix element depth (the type of each individual channel). For example, for a 16-bit signed element array, the method returns ``CV_16S`` . A complete list of matrix types contains the following values:
|
||||
|
||||
* ``CV_8U`` - 8-bit unsigned integers ( ``0..255`` )
|
||||
|
||||
|
||||
@@ -34,7 +34,8 @@ PERF_TEST_P(Size_MatType_ROp, reduceR,
|
||||
declare.in(src, WARMUP_RNG).out(vec);
|
||||
declare.time(100);
|
||||
|
||||
TEST_CYCLE() reduce(src, vec, 0, reduceOp, ddepth);
|
||||
int runs = 15;
|
||||
TEST_CYCLE_MULTIRUN(runs) reduce(src, vec, 0, reduceOp, ddepth);
|
||||
|
||||
SANITY_CHECK(vec, 1);
|
||||
}
|
||||
@@ -65,4 +66,3 @@ PERF_TEST_P(Size_MatType_ROp, reduceC,
|
||||
|
||||
SANITY_CHECK(vec, 1);
|
||||
}
|
||||
|
||||
|
||||
+36
-21
@@ -2850,9 +2850,9 @@ PCA& PCA::operator()(InputArray _data, InputArray __mean, int flags, int maxComp
|
||||
|
||||
if( _mean.data )
|
||||
{
|
||||
CV_Assert( _mean.size() == mean_sz );
|
||||
CV_Assert( _mean.size() == mean_sz );
|
||||
_mean.convertTo(mean, ctype);
|
||||
covar_flags |= CV_COVAR_USE_AVG;
|
||||
covar_flags |= CV_COVAR_USE_AVG;
|
||||
}
|
||||
|
||||
calcCovarMatrix( data, covar, mean, covar_flags, ctype );
|
||||
@@ -2896,6 +2896,36 @@ PCA& PCA::operator()(InputArray _data, InputArray __mean, int flags, int maxComp
|
||||
return *this;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
int computeCumulativeEnergy(const Mat& eigenvalues, double retainedVariance)
|
||||
{
|
||||
CV_DbgAssert( eigenvalues.type() == DataType<T>::type );
|
||||
|
||||
Mat g(eigenvalues.size(), DataType<T>::type);
|
||||
|
||||
for(int ig = 0; ig < g.rows; ig++)
|
||||
{
|
||||
g.at<T>(ig, 0) = 0;
|
||||
for(int im = 0; im <= ig; im++)
|
||||
{
|
||||
g.at<T>(ig,0) += eigenvalues.at<T>(im,0);
|
||||
}
|
||||
}
|
||||
|
||||
int L;
|
||||
|
||||
for(L = 0; L < eigenvalues.rows; L++)
|
||||
{
|
||||
double energy = g.at<T>(L, 0) / g.at<T>(g.rows - 1, 0);
|
||||
if(energy > retainedVariance)
|
||||
break;
|
||||
}
|
||||
|
||||
L = std::max(2, L);
|
||||
|
||||
return L;
|
||||
}
|
||||
|
||||
PCA& PCA::operator()(InputArray _data, InputArray __mean, int flags, double retainedVariance)
|
||||
{
|
||||
Mat data = _data.getMat(), _mean = __mean.getMat();
|
||||
@@ -2972,26 +3002,11 @@ PCA& PCA::operator()(InputArray _data, InputArray __mean, int flags, double reta
|
||||
}
|
||||
|
||||
// compute the cumulative energy content for each eigenvector
|
||||
Mat g(eigenvalues.size(), ctype);
|
||||
|
||||
for(int ig = 0; ig < g.rows; ig++)
|
||||
{
|
||||
g.at<float>(ig,0) = 0;
|
||||
for(int im = 0; im <= ig; im++)
|
||||
{
|
||||
g.at<float>(ig,0) += eigenvalues.at<float>(im,0);
|
||||
}
|
||||
}
|
||||
|
||||
int L;
|
||||
for(L = 0; L < eigenvalues.rows; L++)
|
||||
{
|
||||
double energy = g.at<float>(L, 0) / g.at<float>(g.rows - 1, 0);
|
||||
if(energy > retainedVariance)
|
||||
break;
|
||||
}
|
||||
|
||||
L = std::max(2, L);
|
||||
if (ctype == CV_32F)
|
||||
L = computeCumulativeEnergy<float>(eigenvalues, retainedVariance);
|
||||
else
|
||||
L = computeCumulativeEnergy<double>(eigenvalues, retainedVariance);
|
||||
|
||||
// use clone() to physically copy the data and thus deallocate the original matrices
|
||||
eigenvalues = eigenvalues.rowRange(0,L).clone();
|
||||
|
||||
@@ -189,7 +189,7 @@ For each query descriptor, finds the training descriptors not farther than the s
|
||||
|
||||
:param compactResult: Parameter used when the mask (or masks) is not empty. If ``compactResult`` is false, the ``matches`` vector has the same size as ``queryDescriptors`` rows. If ``compactResult`` is true, the ``matches`` vector does not contain matches for fully masked-out query descriptors.
|
||||
|
||||
:param maxDistance: Threshold for the distance between matched descriptors.
|
||||
:param maxDistance: Threshold for the distance between matched descriptors. Distance means here metric distance (e.g. Hamming distance), not the distance between coordinates (which is measured in Pixels)!
|
||||
|
||||
For each query descriptor, the methods find such training descriptors that the distance between the query descriptor and the training descriptor is equal or smaller than ``maxDistance``. Found matches are returned in the distance increasing order.
|
||||
|
||||
|
||||
@@ -185,7 +185,11 @@ if(HAVE_XIMEA)
|
||||
if(XIMEA_LIBRARY_DIR)
|
||||
link_directories(${XIMEA_LIBRARY_DIR})
|
||||
endif()
|
||||
list(APPEND HIGHGUI_LIBRARIES m3api)
|
||||
if(CMAKE_CL_64)
|
||||
list(APPEND HIGHGUI_LIBRARIES m3apiX64)
|
||||
else()
|
||||
list(APPEND HIGHGUI_LIBRARIES m3api)
|
||||
endif()
|
||||
endif(HAVE_XIMEA)
|
||||
|
||||
if(HAVE_FFMPEG)
|
||||
|
||||
@@ -45,7 +45,16 @@
|
||||
|
||||
#include <unistd.h>
|
||||
#include <stdint.h>
|
||||
#include <sys/select.h>
|
||||
#ifdef WIN32
|
||||
// On Windows, we have no sys/select.h, but we need to pick up
|
||||
// select() which is in winsock2.
|
||||
#ifndef __SYS_SELECT_H__
|
||||
#define __SYS_SELECT_H__ 1
|
||||
#include <winsock2.h>
|
||||
#endif
|
||||
#else
|
||||
#include <sys/select.h>
|
||||
#endif /*WIN32*/
|
||||
#include <dc1394/dc1394.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
|
||||
@@ -3193,8 +3193,10 @@ IplImage* CvCaptureCAM_DShow::retrieveFrame(int)
|
||||
frame = cvCreateImage( cvSize(w,h), 8, 3 );
|
||||
}
|
||||
|
||||
VI.getPixels( index, (uchar*)frame->imageData, false, true );
|
||||
return frame;
|
||||
if (VI.getPixels( index, (uchar*)frame->imageData, false, true ))
|
||||
return frame;
|
||||
else
|
||||
return NULL;
|
||||
}
|
||||
|
||||
double CvCaptureCAM_DShow::getProperty( int property_id )
|
||||
|
||||
@@ -20,25 +20,24 @@ public:
|
||||
virtual IplImage* retrieveFrame(int);
|
||||
virtual int getCaptureDomain() { return CV_CAP_XIAPI; } // Return the type of the capture object: CV_CAP_VFW, etc...
|
||||
|
||||
protected:
|
||||
private:
|
||||
void init();
|
||||
void errMsg(const char* msg, int errNum);
|
||||
void resetCvImage();
|
||||
int getBpp();
|
||||
IplImage* frame;
|
||||
|
||||
HANDLE hmv;
|
||||
DWORD numDevices;
|
||||
XI_IMG image;
|
||||
int width;
|
||||
int height;
|
||||
int format;
|
||||
int timeout;
|
||||
XI_IMG image;
|
||||
};
|
||||
|
||||
/**********************************************************************************/
|
||||
|
||||
CvCapture* cvCreateCameraCapture_XIMEA( int index )
|
||||
{
|
||||
CvCaptureCAM_XIMEA* capture = new CvCaptureCAM_XIMEA;
|
||||
CvCaptureCAM_XIMEA* capture = new CvCaptureCAM_XIMEA;
|
||||
|
||||
if( capture->open( index ))
|
||||
return capture;
|
||||
@@ -79,18 +78,19 @@ bool CvCaptureCAM_XIMEA::open( int wIndex )
|
||||
// always use auto white ballance
|
||||
mvret = xiSetParamInt( hmv, XI_PRM_AUTO_WB, 1);
|
||||
if(mvret != XI_OK) goto error;
|
||||
|
||||
// default image format RGB24
|
||||
mvret = xiSetParamInt( hmv, XI_PRM_IMAGE_DATA_FORMAT, XI_RGB24);
|
||||
if(mvret != XI_OK) goto error;
|
||||
|
||||
int width = 0;
|
||||
mvret = xiGetParamInt( hmv, XI_PRM_WIDTH, &width);
|
||||
if(mvret != XI_OK) goto error;
|
||||
|
||||
int height = 0;
|
||||
mvret = xiGetParamInt( hmv, XI_PRM_HEIGHT, &height);
|
||||
if(mvret != XI_OK) goto error;
|
||||
|
||||
// default image format RGB24
|
||||
format = XI_RGB24;
|
||||
mvret = xiSetParamInt( hmv, XI_PRM_IMAGE_DATA_FORMAT, format);
|
||||
if(mvret != XI_OK) goto error;
|
||||
|
||||
// allocate frame buffer for RGB24 image
|
||||
frame = cvCreateImage(cvSize( width, height), IPL_DEPTH_8U, 3);
|
||||
|
||||
@@ -103,10 +103,10 @@ bool CvCaptureCAM_XIMEA::open( int wIndex )
|
||||
errMsg("StartAcquisition XI_DEVICE failed", mvret);
|
||||
goto error;
|
||||
}
|
||||
|
||||
return true;
|
||||
|
||||
error:
|
||||
errMsg("Open XI_DEVICE failed", mvret);
|
||||
xiCloseDevice(hmv);
|
||||
hmv = NULL;
|
||||
return false;
|
||||
@@ -116,18 +116,19 @@ error:
|
||||
|
||||
void CvCaptureCAM_XIMEA::close()
|
||||
{
|
||||
if(hmv)
|
||||
{
|
||||
xiStopAcquisition(hmv);
|
||||
xiCloseDevice(hmv);
|
||||
hmv = NULL;
|
||||
}
|
||||
if(frame)
|
||||
cvReleaseImage(&frame);
|
||||
|
||||
xiStopAcquisition(hmv);
|
||||
xiCloseDevice(hmv);
|
||||
hmv = NULL;
|
||||
}
|
||||
|
||||
/**********************************************************************************/
|
||||
|
||||
bool CvCaptureCAM_XIMEA::grabFrame()
|
||||
{
|
||||
memset(&image, 0, sizeof(XI_IMG));
|
||||
image.size = sizeof(XI_IMG);
|
||||
int mvret = xiGetImage( hmv, timeout, &image);
|
||||
|
||||
@@ -151,31 +152,18 @@ bool CvCaptureCAM_XIMEA::grabFrame()
|
||||
IplImage* CvCaptureCAM_XIMEA::retrieveFrame(int)
|
||||
{
|
||||
// update cvImage after format has changed
|
||||
if( (int)image.width != width || (int)image.height != height || image.frm != (XI_IMG_FORMAT)format)
|
||||
{
|
||||
cvReleaseImage(&frame);
|
||||
switch( image.frm)
|
||||
{
|
||||
case XI_MONO8 : frame = cvCreateImage(cvSize( image.width, image.height), IPL_DEPTH_8U, 1); break;
|
||||
case XI_MONO16 : frame = cvCreateImage(cvSize( image.width, image.height), IPL_DEPTH_16U, 1); break;
|
||||
case XI_RGB24 : frame = cvCreateImage(cvSize( image.width, image.height), IPL_DEPTH_8U, 3); break;
|
||||
case XI_RGB32 : frame = cvCreateImage(cvSize( image.width, image.height), IPL_DEPTH_8U, 4); break;
|
||||
default :
|
||||
return frame;
|
||||
}
|
||||
// update global image format
|
||||
format = image.frm;
|
||||
width = image.width;
|
||||
height = image.height;
|
||||
}
|
||||
|
||||
resetCvImage();
|
||||
|
||||
// copy pixel data
|
||||
switch( image.frm)
|
||||
{
|
||||
case XI_MONO8 : memcpy( frame->imageData, image.bp, image.width*image.height); break;
|
||||
case XI_MONO16 : memcpy( frame->imageData, image.bp, image.width*image.height*sizeof(WORD)); break;
|
||||
case XI_RGB24 : memcpy( frame->imageData, image.bp, image.width*image.height*3); break;
|
||||
case XI_RGB32 : memcpy( frame->imageData, image.bp, image.width*image.height*sizeof(DWORD)); break;
|
||||
case XI_MONO8 :
|
||||
case XI_RAW8 : memcpy( frame->imageData, image.bp, image.width*image.height); break;
|
||||
case XI_MONO16 :
|
||||
case XI_RAW16 : memcpy( frame->imageData, image.bp, image.width*image.height*sizeof(WORD)); break;
|
||||
case XI_RGB24 :
|
||||
case XI_RGB_PLANAR : memcpy( frame->imageData, image.bp, image.width*image.height*3); break;
|
||||
case XI_RGB32 : memcpy( frame->imageData, image.bp, image.width*image.height*4); break;
|
||||
default: break;
|
||||
}
|
||||
return frame;
|
||||
@@ -183,6 +171,35 @@ IplImage* CvCaptureCAM_XIMEA::retrieveFrame(int)
|
||||
|
||||
/**********************************************************************************/
|
||||
|
||||
void CvCaptureCAM_XIMEA::resetCvImage()
|
||||
{
|
||||
int width = 0, height = 0, format = 0;
|
||||
xiGetParamInt( hmv, XI_PRM_WIDTH, &width);
|
||||
xiGetParamInt( hmv, XI_PRM_HEIGHT, &height);
|
||||
xiGetParamInt( hmv, XI_PRM_IMAGE_DATA_FORMAT, &format);
|
||||
|
||||
if( (int)image.width != width || (int)image.height != height || image.frm != (XI_IMG_FORMAT)format)
|
||||
{
|
||||
if(frame) cvReleaseImage(&frame);
|
||||
frame = NULL;
|
||||
|
||||
switch( image.frm)
|
||||
{
|
||||
case XI_MONO8 :
|
||||
case XI_RAW8 : frame = cvCreateImage(cvSize( image.width, image.height), IPL_DEPTH_8U, 1); break;
|
||||
case XI_MONO16 :
|
||||
case XI_RAW16 : frame = cvCreateImage(cvSize( image.width, image.height), IPL_DEPTH_16U, 1); break;
|
||||
case XI_RGB24 :
|
||||
case XI_RGB_PLANAR : frame = cvCreateImage(cvSize( image.width, image.height), IPL_DEPTH_8U, 3); break;
|
||||
case XI_RGB32 : frame = cvCreateImage(cvSize( image.width, image.height), IPL_DEPTH_8U, 4); break;
|
||||
default :
|
||||
return;
|
||||
}
|
||||
}
|
||||
cvZero(frame);
|
||||
}
|
||||
/**********************************************************************************/
|
||||
|
||||
double CvCaptureCAM_XIMEA::getProperty( int property_id )
|
||||
{
|
||||
if(hmv == NULL)
|
||||
@@ -238,20 +255,14 @@ bool CvCaptureCAM_XIMEA::setProperty( int property_id, double value )
|
||||
switch(property_id)
|
||||
{
|
||||
// OCV parameters
|
||||
case CV_CAP_PROP_FRAME_WIDTH : mvret = xiSetParamInt( hmv, XI_PRM_WIDTH, ival);
|
||||
if(mvret == XI_OK) width = ival;
|
||||
break;
|
||||
case CV_CAP_PROP_FRAME_HEIGHT : mvret = xiSetParamInt( hmv, XI_PRM_HEIGHT, ival);
|
||||
if(mvret == XI_OK) height = ival;
|
||||
break;
|
||||
case CV_CAP_PROP_FRAME_WIDTH : mvret = xiSetParamInt( hmv, XI_PRM_WIDTH, ival); break;
|
||||
case CV_CAP_PROP_FRAME_HEIGHT : mvret = xiSetParamInt( hmv, XI_PRM_HEIGHT, ival); break;
|
||||
case CV_CAP_PROP_FPS : mvret = xiSetParamFloat( hmv, XI_PRM_FRAMERATE, fval); break;
|
||||
case CV_CAP_PROP_GAIN : mvret = xiSetParamFloat( hmv, XI_PRM_GAIN, fval); break;
|
||||
case CV_CAP_PROP_EXPOSURE : mvret = xiSetParamInt( hmv, XI_PRM_EXPOSURE, ival); break;
|
||||
// XIMEA camera properties
|
||||
case CV_CAP_PROP_XI_DOWNSAMPLING : mvret = xiSetParamInt( hmv, XI_PRM_DOWNSAMPLING, ival); break;
|
||||
case CV_CAP_PROP_XI_DATA_FORMAT : mvret = xiSetParamInt( hmv, XI_PRM_IMAGE_DATA_FORMAT, ival);
|
||||
if(mvret == XI_OK) format = ival;
|
||||
break;
|
||||
case CV_CAP_PROP_XI_DATA_FORMAT : mvret = xiSetParamInt( hmv, XI_PRM_IMAGE_DATA_FORMAT, ival); break;
|
||||
case CV_CAP_PROP_XI_OFFSET_X : mvret = xiSetParamInt( hmv, XI_PRM_OFFSET_X, ival); break;
|
||||
case CV_CAP_PROP_XI_OFFSET_Y : mvret = xiSetParamInt( hmv, XI_PRM_OFFSET_Y, ival); break;
|
||||
case CV_CAP_PROP_XI_TRG_SOURCE : mvret = xiSetParamInt( hmv, XI_PRM_TRG_SOURCE, ival); break;
|
||||
@@ -288,7 +299,7 @@ bool CvCaptureCAM_XIMEA::setProperty( int property_id, double value )
|
||||
void CvCaptureCAM_XIMEA::errMsg(const char* msg, int errNum)
|
||||
{
|
||||
#if defined WIN32 || defined _WIN32
|
||||
char buf[512];
|
||||
char buf[512]="";
|
||||
sprintf( buf, "%s : %d\n", msg, errNum);
|
||||
OutputDebugString(buf);
|
||||
#else
|
||||
@@ -296,4 +307,22 @@ void CvCaptureCAM_XIMEA::errMsg(const char* msg, int errNum)
|
||||
#endif
|
||||
}
|
||||
|
||||
/**********************************************************************************/
|
||||
|
||||
int CvCaptureCAM_XIMEA::getBpp()
|
||||
{
|
||||
switch( image.frm)
|
||||
{
|
||||
case XI_MONO8 :
|
||||
case XI_RAW8 : return 1;
|
||||
case XI_MONO16 :
|
||||
case XI_RAW16 : return 2;
|
||||
case XI_RGB24 :
|
||||
case XI_RGB_PLANAR : return 3;
|
||||
case XI_RGB32 : return 4;
|
||||
default :
|
||||
return 0;
|
||||
}
|
||||
}
|
||||
|
||||
/**********************************************************************************/
|
||||
@@ -256,12 +256,17 @@ namespace
|
||||
|
||||
void cv::imshow( const String& winname, InputArray _img )
|
||||
{
|
||||
const Size size = _img.size();
|
||||
#ifndef HAVE_OPENGL
|
||||
Mat img = _img.getMat();
|
||||
CvMat c_img = img;
|
||||
cvShowImage(winname.c_str(), &c_img);
|
||||
CV_Assert(size.width>0 && size.height>0);
|
||||
{
|
||||
Mat img = _img.getMat();
|
||||
CvMat c_img = img;
|
||||
cvShowImage(winname.c_str(), &c_img);
|
||||
}
|
||||
#else
|
||||
const double useGl = getWindowProperty(winname, WND_PROP_OPENGL);
|
||||
CV_Assert(size.width>0 && size.height>0);
|
||||
|
||||
if (useGl <= 0)
|
||||
{
|
||||
@@ -275,7 +280,6 @@ void cv::imshow( const String& winname, InputArray _img )
|
||||
|
||||
if (autoSize > 0)
|
||||
{
|
||||
Size size = _img.size();
|
||||
resizeWindow(winname, size.width, size.height);
|
||||
}
|
||||
|
||||
|
||||
@@ -258,7 +258,8 @@ PERF_TEST_P(Size_CvtMode, cvtColor8u,
|
||||
declare.time(100);
|
||||
declare.in(src, WARMUP_RNG).out(dst);
|
||||
|
||||
TEST_CYCLE() cvtColor(src, dst, mode, ch.dcn);
|
||||
int runs = sz.width <= 320 ? 70 : 1;
|
||||
TEST_CYCLE_MULTIRUN(runs) cvtColor(src, dst, mode, ch.dcn);
|
||||
|
||||
SANITY_CHECK(dst, 1);
|
||||
}
|
||||
@@ -334,7 +335,8 @@ PERF_TEST_P(Size_CvtMode3, cvtColorRGB2YUV420p,
|
||||
declare.time(100);
|
||||
declare.in(src, WARMUP_RNG).out(dst);
|
||||
|
||||
TEST_CYCLE() cvtColor(src, dst, mode, ch.dcn);
|
||||
int runs = (sz.width <= 640) ? 10 : 1;
|
||||
TEST_CYCLE_MULTIRUN(runs) cvtColor(src, dst, mode, ch.dcn);
|
||||
|
||||
SANITY_CHECK(dst, 1);
|
||||
}
|
||||
|
||||
@@ -19,7 +19,8 @@ PERF_TEST_P(Size_MatType, erode, TYPICAL_MATS_MORPH)
|
||||
|
||||
declare.in(src, WARMUP_RNG).out(dst);
|
||||
|
||||
TEST_CYCLE() erode(src, dst, noArray());
|
||||
int runs = (sz.width <= 320) ? 15 : 1;
|
||||
TEST_CYCLE_MULTIRUN(runs) erode(src, dst, noArray());
|
||||
|
||||
SANITY_CHECK(dst);
|
||||
}
|
||||
|
||||
@@ -63,7 +63,8 @@ PERF_TEST_P( TestRemap, Remap,
|
||||
|
||||
declare.in(src, WARMUP_RNG).out(dst).time(20);
|
||||
|
||||
TEST_CYCLE() remap(src, dst, map1, map2, inter_type);
|
||||
int runs = (sz.width <= 640) ? 3 : 1;
|
||||
TEST_CYCLE_MULTIRUN(runs) remap(src, dst, map1, map2, inter_type);
|
||||
|
||||
SANITY_CHECK(dst);
|
||||
}
|
||||
|
||||
@@ -32,7 +32,7 @@ PERF_TEST_P(Size_MatType_ThreshType, threshold,
|
||||
|
||||
declare.in(src, WARMUP_RNG).out(dst);
|
||||
|
||||
int runs = (sz.width <= 640) ? 8 : 1;
|
||||
int runs = (sz.width <= 640) ? 40 : 1;
|
||||
TEST_CYCLE_MULTIRUN(runs) threshold(src, dst, thresh, maxval, threshType);
|
||||
|
||||
SANITY_CHECK(dst);
|
||||
|
||||
@@ -80,10 +80,10 @@ public abstract class CameraBridgeViewBase extends SurfaceView implements Surfac
|
||||
mMaxHeight = MAX_UNSPECIFIED;
|
||||
styledAttrs.recycle();
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Sets the camera index
|
||||
* @param camera index
|
||||
* @param cameraIndex new camera index
|
||||
*/
|
||||
public void setCameraIndex(int cameraIndex) {
|
||||
this.mCameraIndex = cameraIndex;
|
||||
|
||||
@@ -14,7 +14,7 @@ public class MatOfByte extends Mat {
|
||||
|
||||
protected MatOfByte(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -25,7 +25,7 @@ public class MatOfByte extends Mat {
|
||||
|
||||
public MatOfByte(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -14,7 +14,7 @@ public class MatOfDouble extends Mat {
|
||||
|
||||
protected MatOfDouble(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -25,7 +25,7 @@ public class MatOfDouble extends Mat {
|
||||
|
||||
public MatOfDouble(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -14,7 +14,7 @@ public class MatOfFloat extends Mat {
|
||||
|
||||
protected MatOfFloat(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -25,7 +25,7 @@ public class MatOfFloat extends Mat {
|
||||
|
||||
public MatOfFloat(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -14,7 +14,7 @@ public class MatOfFloat4 extends Mat {
|
||||
|
||||
protected MatOfFloat4(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -25,7 +25,7 @@ public class MatOfFloat4 extends Mat {
|
||||
|
||||
public MatOfFloat4(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -14,7 +14,7 @@ public class MatOfFloat6 extends Mat {
|
||||
|
||||
protected MatOfFloat6(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -25,7 +25,7 @@ public class MatOfFloat6 extends Mat {
|
||||
|
||||
public MatOfFloat6(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -15,7 +15,7 @@ public class MatOfInt extends Mat {
|
||||
|
||||
protected MatOfInt(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -26,7 +26,7 @@ public class MatOfInt extends Mat {
|
||||
|
||||
public MatOfInt(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -15,7 +15,7 @@ public class MatOfInt4 extends Mat {
|
||||
|
||||
protected MatOfInt4(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -26,7 +26,7 @@ public class MatOfInt4 extends Mat {
|
||||
|
||||
public MatOfInt4(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -16,7 +16,7 @@ public class MatOfKeyPoint extends Mat {
|
||||
|
||||
protected MatOfKeyPoint(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -27,7 +27,7 @@ public class MatOfKeyPoint extends Mat {
|
||||
|
||||
public MatOfKeyPoint(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -14,7 +14,7 @@ public class MatOfPoint extends Mat {
|
||||
|
||||
protected MatOfPoint(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -25,7 +25,7 @@ public class MatOfPoint extends Mat {
|
||||
|
||||
public MatOfPoint(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -14,7 +14,7 @@ public class MatOfPoint2f extends Mat {
|
||||
|
||||
protected MatOfPoint2f(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -25,7 +25,7 @@ public class MatOfPoint2f extends Mat {
|
||||
|
||||
public MatOfPoint2f(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -14,7 +14,7 @@ public class MatOfPoint3 extends Mat {
|
||||
|
||||
protected MatOfPoint3(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -25,7 +25,7 @@ public class MatOfPoint3 extends Mat {
|
||||
|
||||
public MatOfPoint3(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -14,7 +14,7 @@ public class MatOfPoint3f extends Mat {
|
||||
|
||||
protected MatOfPoint3f(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -25,7 +25,7 @@ public class MatOfPoint3f extends Mat {
|
||||
|
||||
public MatOfPoint3f(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -15,7 +15,7 @@ public class MatOfRect extends Mat {
|
||||
|
||||
protected MatOfRect(long addr) {
|
||||
super(addr);
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
@@ -26,7 +26,7 @@ public class MatOfRect extends Mat {
|
||||
|
||||
public MatOfRect(Mat m) {
|
||||
super(m, Range.all());
|
||||
if(checkVector(_channels, _depth) < 0 )
|
||||
if( !empty() && checkVector(_channels, _depth) < 0 )
|
||||
throw new IllegalArgumentException("Incomatible Mat");
|
||||
//FIXME: do we need release() here?
|
||||
}
|
||||
|
||||
@@ -251,7 +251,7 @@ void CvANN_MLP::create( const CvMat* _layer_sizes, int _activ_func,
|
||||
buf_sz += (l_dst[0] + l_dst[l_count-1]*2)*2;
|
||||
|
||||
CV_CALL( wbuf = cvCreateMat( 1, buf_sz, CV_64F ));
|
||||
CV_CALL( weights = (double**)cvAlloc( (l_count+1)*sizeof(weights[0]) ));
|
||||
CV_CALL( weights = (double**)cvAlloc( (l_count+2)*sizeof(weights[0]) ));
|
||||
|
||||
weights[0] = wbuf->data.db;
|
||||
weights[1] = weights[0] + l_dst[0]*2;
|
||||
|
||||
@@ -774,9 +774,6 @@ void SIFT::operator()(InputArray _image, InputArray _mask,
|
||||
findScaleSpaceExtrema(gpyr, dogpyr, keypoints);
|
||||
KeyPointsFilter::removeDuplicated( keypoints );
|
||||
|
||||
if( !mask.empty() )
|
||||
KeyPointsFilter::runByPixelsMask( keypoints, mask );
|
||||
|
||||
if( nfeatures > 0 )
|
||||
KeyPointsFilter::retainBest(keypoints, nfeatures);
|
||||
//t = (double)getTickCount() - t;
|
||||
@@ -791,6 +788,9 @@ void SIFT::operator()(InputArray _image, InputArray _mask,
|
||||
kpt.pt *= scale;
|
||||
kpt.size *= scale;
|
||||
}
|
||||
|
||||
if( !mask.empty() )
|
||||
KeyPointsFilter::runByPixelsMask( keypoints, mask );
|
||||
}
|
||||
else
|
||||
{
|
||||
|
||||
@@ -117,9 +117,6 @@ namespace cv
|
||||
//the devnum is the index of the selected device in DeviceName vector of INfo
|
||||
CV_EXPORTS void setDevice(Info &oclinfo, int devnum = 0);
|
||||
|
||||
//optional function, if you want save opencl binary kernel to the file, set its path
|
||||
CV_EXPORTS void setBinpath(const char *path);
|
||||
|
||||
//The two functions below enable other opencl program to use ocl module's cl_context and cl_command_queue
|
||||
//returns cl_context *
|
||||
CV_EXPORTS void* getoclContext();
|
||||
@@ -133,6 +130,9 @@ namespace cv
|
||||
//getDevice also need to be called before this function
|
||||
CV_EXPORTS void setDeviceEx(Info &oclinfo, void *ctx, void *qu, int devnum = 0);
|
||||
|
||||
//returns true when global OpenCL context is initialized
|
||||
CV_EXPORTS bool initialized();
|
||||
|
||||
//////////////////////////////// OpenCL context ////////////////////////
|
||||
//This is a global singleton class used to represent a OpenCL context.
|
||||
class CV_EXPORTS Context
|
||||
@@ -140,7 +140,7 @@ namespace cv
|
||||
protected:
|
||||
Context();
|
||||
friend class std::auto_ptr<Context>;
|
||||
|
||||
friend bool initialized();
|
||||
private:
|
||||
static std::auto_ptr<Context> clCxt;
|
||||
static int val;
|
||||
@@ -178,6 +178,29 @@ namespace cv
|
||||
bool finish = true, bool measureKernelTime = false,
|
||||
bool cleanUp = true);
|
||||
|
||||
//! Enable or disable OpenCL program binary caching onto local disk
|
||||
// After a program (*.cl files in opencl/ folder) is built at runtime, we allow the
|
||||
// compiled OpenCL program to be cached to the path automatically as "path/*.clb"
|
||||
// binary file, which will be reused when the OpenCV executable is started again.
|
||||
//
|
||||
// Caching mode is controlled by the following enums
|
||||
// Notes
|
||||
// 1. the feature is by default enabled when OpenCV is built in release mode.
|
||||
// 2. the CACHE_DEBUG / CACHE_RELEASE flags only effectively work with MSVC compiler;
|
||||
// for GNU compilers, the function always treats the build as release mode (enabled by default).
|
||||
enum
|
||||
{
|
||||
CACHE_NONE = 0, // do not cache OpenCL binary
|
||||
CACHE_DEBUG = 0x1 << 0, // cache OpenCL binary when built in debug mode (only work with MSVC)
|
||||
CACHE_RELEASE = 0x1 << 1, // default behavior, only cache when built in release mode (only work with MSVC)
|
||||
CACHE_ALL = CACHE_DEBUG | CACHE_RELEASE, // always cache opencl binary
|
||||
CACHE_UPDATE = 0x1 << 2 // if the binary cache file with the same name is already on the disk, it will be updated.
|
||||
};
|
||||
CV_EXPORTS void setBinaryDiskCache(int mode = CACHE_RELEASE, cv::String path = "./");
|
||||
|
||||
//! set where binary cache to be saved to
|
||||
CV_EXPORTS void setBinpath(const char *path);
|
||||
|
||||
class CV_EXPORTS oclMatExpr;
|
||||
//////////////////////////////// oclMat ////////////////////////////////
|
||||
class CV_EXPORTS oclMat
|
||||
@@ -482,6 +505,25 @@ namespace cv
|
||||
CV_EXPORTS void calcHist(const oclMat &mat_src, oclMat &mat_hist);
|
||||
//! only 8UC1 and 256 bins is supported now
|
||||
CV_EXPORTS void equalizeHist(const oclMat &mat_src, oclMat &mat_dst);
|
||||
|
||||
//! only 8UC1 is supported now
|
||||
class CV_EXPORTS CLAHE
|
||||
{
|
||||
public:
|
||||
virtual void apply(const oclMat &src, oclMat &dst) = 0;
|
||||
|
||||
virtual void setClipLimit(double clipLimit) = 0;
|
||||
virtual double getClipLimit() const = 0;
|
||||
|
||||
virtual void setTilesGridSize(Size tileGridSize) = 0;
|
||||
virtual Size getTilesGridSize() const = 0;
|
||||
|
||||
virtual void collectGarbage() = 0;
|
||||
|
||||
virtual ~CLAHE() {}
|
||||
};
|
||||
CV_EXPORTS Ptr<cv::ocl::CLAHE> createCLAHE(double clipLimit = 40.0, Size tileGridSize = Size(8, 8));
|
||||
|
||||
//! bilateralFilter
|
||||
// supports 8UC1 8UC4
|
||||
CV_EXPORTS void bilateralFilter(const oclMat& src, oclMat& dst, int d, double sigmaColor, double sigmaSpave, int borderType=BORDER_DEFAULT);
|
||||
|
||||
@@ -922,3 +922,50 @@ PERFTEST(remap)
|
||||
|
||||
}
|
||||
}
|
||||
///////////// CLAHE ////////////////////////
|
||||
PERFTEST(CLAHE)
|
||||
{
|
||||
Mat src, dst, ocl_dst;
|
||||
cv::ocl::oclMat d_src, d_dst;
|
||||
int all_type[] = {CV_8UC1};
|
||||
std::string type_name[] = {"CV_8UC1"};
|
||||
|
||||
double clipLimit = 40.0;
|
||||
|
||||
cv::Ptr<cv::CLAHE> clahe = cv::createCLAHE(clipLimit);
|
||||
cv::Ptr<cv::ocl::CLAHE> d_clahe = cv::ocl::createCLAHE(clipLimit);
|
||||
|
||||
for (int size = Min_Size; size <= Max_Size; size *= Multiple)
|
||||
{
|
||||
for (size_t j = 0; j < sizeof(all_type) / sizeof(int); j++)
|
||||
{
|
||||
SUBTEST << size << 'x' << size << "; " << type_name[j] ;
|
||||
|
||||
gen(src, size, size, all_type[j], 0, 256);
|
||||
|
||||
CPU_ON;
|
||||
clahe->apply(src, dst);
|
||||
CPU_OFF;
|
||||
|
||||
d_src.upload(src);
|
||||
|
||||
WARMUP_ON;
|
||||
d_clahe->apply(d_src, d_dst);
|
||||
WARMUP_OFF;
|
||||
|
||||
ocl_dst = d_dst;
|
||||
|
||||
TestSystem::instance().ExpectedMatNear(dst, ocl_dst, 1.0);
|
||||
|
||||
GPU_ON;
|
||||
d_clahe->apply(d_src, d_dst);
|
||||
GPU_OFF;
|
||||
|
||||
GPU_FULL_ON;
|
||||
d_src.upload(src);
|
||||
d_clahe->apply(d_src, d_dst);
|
||||
d_dst.download(dst);
|
||||
GPU_FULL_OFF;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -26,6 +26,7 @@
|
||||
// Wu Zailong, bullet@yeah.net
|
||||
// Wenju He, wenju@multicorewareinc.com
|
||||
// Peng Xiao, pengxiao@outlook.com
|
||||
// Sen Liu, swjtuls1987@126.com
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
@@ -81,6 +82,7 @@ namespace cv
|
||||
extern const char *imgproc_calcMinEigenVal;
|
||||
extern const char *imgproc_convolve;
|
||||
extern const char *imgproc_mulAndScaleSpectrums;
|
||||
extern const char *imgproc_clahe;
|
||||
////////////////////////////////////OpenCL call wrappers////////////////////////////
|
||||
|
||||
template <typename T> struct index_and_sizeof;
|
||||
@@ -1505,6 +1507,189 @@ namespace cv
|
||||
openCLExecuteKernel(clCxt, &imgproc_histogram, kernelName, globalThreads, localThreads, args, -1, -1);
|
||||
LUT(mat_src, lut, mat_dst);
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////
|
||||
// CLAHE
|
||||
namespace clahe
|
||||
{
|
||||
inline int divUp(int total, int grain)
|
||||
{
|
||||
return (total + grain - 1) / grain * grain;
|
||||
}
|
||||
|
||||
static void calcLut(const oclMat &src, oclMat &dst,
|
||||
const int tilesX, const int tilesY, const cv::Size tileSize,
|
||||
const int clipLimit, const float lutScale)
|
||||
{
|
||||
cl_int2 tile_size;
|
||||
tile_size.s[0] = tileSize.width;
|
||||
tile_size.s[1] = tileSize.height;
|
||||
|
||||
std::vector<std::pair<size_t , const void *> > args;
|
||||
args.push_back( std::make_pair( sizeof(cl_mem), (void *)&src.data ));
|
||||
args.push_back( std::make_pair( sizeof(cl_mem), (void *)&dst.data ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.step ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst.step ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int2), (void *)&tile_size ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&tilesX ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&clipLimit ));
|
||||
args.push_back( std::make_pair( sizeof(cl_float), (void *)&lutScale ));
|
||||
|
||||
String kernelName = "calcLut";
|
||||
size_t localThreads[3] = { 32, 8, 1 };
|
||||
size_t globalThreads[3] = { tilesX * localThreads[0], tilesY * localThreads[1], 1 };
|
||||
bool is_cpu = queryDeviceInfo<IS_CPU_DEVICE, bool>();
|
||||
if (is_cpu)
|
||||
{
|
||||
openCLExecuteKernel(Context::getContext(), &imgproc_clahe, kernelName, globalThreads, localThreads, args, -1, -1, (char*)" -D CPU");
|
||||
}
|
||||
else
|
||||
{
|
||||
cl_kernel kernel = openCLGetKernelFromSource(Context::getContext(), &imgproc_clahe, kernelName);
|
||||
int wave_size = queryDeviceInfo<WAVEFRONT_SIZE, int>(kernel);
|
||||
openCLSafeCall(clReleaseKernel(kernel));
|
||||
|
||||
static char opt[20] = {0};
|
||||
sprintf(opt, " -D WAVE_SIZE=%d", wave_size);
|
||||
openCLExecuteKernel(Context::getContext(), &imgproc_clahe, kernelName, globalThreads, localThreads, args, -1, -1, opt);
|
||||
}
|
||||
}
|
||||
|
||||
static void transform(const oclMat &src, oclMat &dst, const oclMat &lut,
|
||||
const int tilesX, const int tilesY, const cv::Size tileSize)
|
||||
{
|
||||
cl_int2 tile_size;
|
||||
tile_size.s[0] = tileSize.width;
|
||||
tile_size.s[1] = tileSize.height;
|
||||
|
||||
std::vector<std::pair<size_t , const void *> > args;
|
||||
args.push_back( std::make_pair( sizeof(cl_mem), (void *)&src.data ));
|
||||
args.push_back( std::make_pair( sizeof(cl_mem), (void *)&dst.data ));
|
||||
args.push_back( std::make_pair( sizeof(cl_mem), (void *)&lut.data ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.step ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst.step ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&lut.step ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.cols ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.rows ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int2), (void *)&tile_size ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&tilesX ));
|
||||
args.push_back( std::make_pair( sizeof(cl_int), (void *)&tilesY ));
|
||||
|
||||
String kernelName = "transform";
|
||||
size_t localThreads[3] = { 32, 8, 1 };
|
||||
size_t globalThreads[3] = { divUp(src.cols, localThreads[0]), divUp(src.rows, localThreads[1]), 1 };
|
||||
|
||||
openCLExecuteKernel(Context::getContext(), &imgproc_clahe, kernelName, globalThreads, localThreads, args, -1, -1);
|
||||
}
|
||||
}
|
||||
|
||||
namespace
|
||||
{
|
||||
class CLAHE_Impl : public cv::ocl::CLAHE
|
||||
{
|
||||
public:
|
||||
CLAHE_Impl(double clipLimit = 40.0, int tilesX = 8, int tilesY = 8);
|
||||
|
||||
cv::AlgorithmInfo* info() const;
|
||||
|
||||
void apply(const oclMat &src, oclMat &dst);
|
||||
|
||||
void setClipLimit(double clipLimit);
|
||||
double getClipLimit() const;
|
||||
|
||||
void setTilesGridSize(cv::Size tileGridSize);
|
||||
cv::Size getTilesGridSize() const;
|
||||
|
||||
void collectGarbage();
|
||||
|
||||
private:
|
||||
double clipLimit_;
|
||||
int tilesX_;
|
||||
int tilesY_;
|
||||
|
||||
oclMat srcExt_;
|
||||
oclMat lut_;
|
||||
};
|
||||
|
||||
CLAHE_Impl::CLAHE_Impl(double clipLimit, int tilesX, int tilesY) :
|
||||
clipLimit_(clipLimit), tilesX_(tilesX), tilesY_(tilesY)
|
||||
{
|
||||
}
|
||||
|
||||
void CLAHE_Impl::apply(const oclMat &src, oclMat &dst)
|
||||
{
|
||||
CV_Assert( src.type() == CV_8UC1 );
|
||||
|
||||
dst.create( src.size(), src.type() );
|
||||
|
||||
const int histSize = 256;
|
||||
|
||||
ensureSizeIsEnough(tilesX_ * tilesY_, histSize, CV_8UC1, lut_);
|
||||
|
||||
cv::Size tileSize;
|
||||
oclMat srcForLut;
|
||||
|
||||
if (src.cols % tilesX_ == 0 && src.rows % tilesY_ == 0)
|
||||
{
|
||||
tileSize = cv::Size(src.cols / tilesX_, src.rows / tilesY_);
|
||||
srcForLut = src;
|
||||
}
|
||||
else
|
||||
{
|
||||
cv::ocl::copyMakeBorder(src, srcExt_, 0, tilesY_ - (src.rows % tilesY_), 0, tilesX_ - (src.cols % tilesX_), cv::BORDER_REFLECT_101, cv::Scalar());
|
||||
|
||||
tileSize = cv::Size(srcExt_.cols / tilesX_, srcExt_.rows / tilesY_);
|
||||
srcForLut = srcExt_;
|
||||
}
|
||||
|
||||
const int tileSizeTotal = tileSize.area();
|
||||
const float lutScale = static_cast<float>(histSize - 1) / tileSizeTotal;
|
||||
|
||||
int clipLimit = 0;
|
||||
if (clipLimit_ > 0.0)
|
||||
{
|
||||
clipLimit = static_cast<int>(clipLimit_ * tileSizeTotal / histSize);
|
||||
clipLimit = std::max(clipLimit, 1);
|
||||
}
|
||||
|
||||
clahe::calcLut(srcForLut, lut_, tilesX_, tilesY_, tileSize, clipLimit, lutScale);
|
||||
//finish();
|
||||
clahe::transform(src, dst, lut_, tilesX_, tilesY_, tileSize);
|
||||
}
|
||||
|
||||
void CLAHE_Impl::setClipLimit(double clipLimit)
|
||||
{
|
||||
clipLimit_ = clipLimit;
|
||||
}
|
||||
|
||||
double CLAHE_Impl::getClipLimit() const
|
||||
{
|
||||
return clipLimit_;
|
||||
}
|
||||
|
||||
void CLAHE_Impl::setTilesGridSize(cv::Size tileGridSize)
|
||||
{
|
||||
tilesX_ = tileGridSize.width;
|
||||
tilesY_ = tileGridSize.height;
|
||||
}
|
||||
|
||||
cv::Size CLAHE_Impl::getTilesGridSize() const
|
||||
{
|
||||
return cv::Size(tilesX_, tilesY_);
|
||||
}
|
||||
|
||||
void CLAHE_Impl::collectGarbage()
|
||||
{
|
||||
srcExt_.release();
|
||||
lut_.release();
|
||||
}
|
||||
}
|
||||
|
||||
cv::Ptr<cv::ocl::CLAHE> createCLAHE(double clipLimit, cv::Size tileGridSize)
|
||||
{
|
||||
return new CLAHE_Impl(clipLimit, tileGridSize.width, tileGridSize.height);
|
||||
}
|
||||
|
||||
//////////////////////////////////bilateralFilter////////////////////////////////////////////////////
|
||||
static void
|
||||
oclbilateralFilter_8u( const oclMat &src, oclMat &dst, int d,
|
||||
|
||||
@@ -121,8 +121,9 @@ namespace cv
|
||||
cacheSize = 0;
|
||||
}
|
||||
|
||||
|
||||
struct Info::Impl
|
||||
// not to be exported to dynamic lib
|
||||
void setBinaryDiskCacheImpl(int mode, String path, Info::Impl * impl);
|
||||
struct Info::Impl
|
||||
{
|
||||
cl_platform_id oclplatform;
|
||||
std::vector<cl_device_id> devices;
|
||||
@@ -140,22 +141,12 @@ namespace cv
|
||||
char extra_options[512];
|
||||
int double_support;
|
||||
int unified_memory; //1 means integrated GPU, otherwise this value is 0
|
||||
bool enable_disk_cache;
|
||||
bool update_disk_cache;
|
||||
String binpath;
|
||||
int refcounter;
|
||||
|
||||
Impl()
|
||||
{
|
||||
refcounter = 1;
|
||||
oclplatform = 0;
|
||||
oclcontext = 0;
|
||||
clCmdQueue = 0;
|
||||
devnum = -1;
|
||||
maxComputeUnits = 0;
|
||||
maxWorkGroupSize = 0;
|
||||
memset(extra_options, 0, 512);
|
||||
double_support = 0;
|
||||
unified_memory = 0;
|
||||
}
|
||||
Impl();
|
||||
|
||||
void setDevice(void *ctx, void *q, int devnum);
|
||||
|
||||
@@ -180,6 +171,25 @@ namespace cv
|
||||
void releaseResources();
|
||||
};
|
||||
|
||||
Info::Impl::Impl()
|
||||
:oclplatform(0),
|
||||
oclcontext(0),
|
||||
clCmdQueue(0),
|
||||
devnum(-1),
|
||||
maxWorkGroupSize(0),
|
||||
maxDimensions(0),
|
||||
maxComputeUnits(0),
|
||||
double_support(0),
|
||||
unified_memory(0),
|
||||
enable_disk_cache(false),
|
||||
update_disk_cache(false),
|
||||
binpath("./"),
|
||||
refcounter(1)
|
||||
{
|
||||
memset(extra_options, 0, 512);
|
||||
setBinaryDiskCacheImpl(CACHE_RELEASE, String("./"), this);
|
||||
}
|
||||
|
||||
void Info::Impl::releaseResources()
|
||||
{
|
||||
devnum = -1;
|
||||
@@ -498,6 +508,24 @@ namespace cv
|
||||
return openCLGetKernelFromSource(clCxt, source, kernelName, NULL);
|
||||
}
|
||||
|
||||
void setBinaryDiskCacheImpl(int mode, String path, Info::Impl * impl)
|
||||
{
|
||||
impl->update_disk_cache = (mode & CACHE_UPDATE) == CACHE_UPDATE;
|
||||
impl->enable_disk_cache =
|
||||
#ifdef _DEBUG
|
||||
(mode & CACHE_DEBUG) == CACHE_DEBUG;
|
||||
#else
|
||||
(mode & CACHE_RELEASE) == CACHE_RELEASE;
|
||||
#endif
|
||||
if(impl->enable_disk_cache && !path.empty())
|
||||
{
|
||||
impl->binpath = path;
|
||||
}
|
||||
}
|
||||
void setBinaryDiskCache(int mode, cv::String path)
|
||||
{
|
||||
setBinaryDiskCacheImpl(mode, path, Context::getContext()->impl);
|
||||
}
|
||||
|
||||
void setBinpath(const char *path)
|
||||
{
|
||||
@@ -577,8 +605,8 @@ namespace cv
|
||||
filename = clCxt->impl->binpath + kernelName + "_" + clCxt->impl->devName[clCxt->impl->devnum] + ".clb";
|
||||
}
|
||||
|
||||
FILE *fp = fopen(filename.c_str(), "rb");
|
||||
if(fp == NULL || clCxt->impl->binpath.size() == 0) //we should generate a binary file for the first time.
|
||||
FILE *fp = clCxt->impl->enable_disk_cache ? fopen(filename.c_str(), "rb") : NULL;
|
||||
if(fp == NULL || clCxt->impl->update_disk_cache)
|
||||
{
|
||||
if(fp != NULL)
|
||||
fclose(fp);
|
||||
@@ -587,7 +615,7 @@ namespace cv
|
||||
clCxt->impl->oclcontext, 1, source, NULL, &status);
|
||||
openCLVerifyCall(status);
|
||||
status = clBuildProgram(program, 1, &(clCxt->impl->devices[clCxt->impl->devnum]), all_build_options, NULL, NULL);
|
||||
if(status == CL_SUCCESS && clCxt->impl->binpath.size())
|
||||
if(status == CL_SUCCESS && clCxt->impl->enable_disk_cache)
|
||||
savetofile(clCxt, program, filename.c_str());
|
||||
}
|
||||
else
|
||||
@@ -921,6 +949,14 @@ namespace cv
|
||||
int Context::val = 0;
|
||||
static Mutex cs;
|
||||
static volatile int context_tear_down = 0;
|
||||
|
||||
bool initialized()
|
||||
{
|
||||
return *((volatile int*)&Context::val) != 0 &&
|
||||
Context::clCxt->impl->clCmdQueue != NULL&&
|
||||
Context::clCxt->impl->oclcontext != NULL;
|
||||
}
|
||||
|
||||
Context* Context::getContext()
|
||||
{
|
||||
if(*((volatile int*)&val) != 1)
|
||||
@@ -934,8 +970,6 @@ namespace cv
|
||||
clCxt.reset(new Context);
|
||||
std::vector<Info> oclinfo;
|
||||
CV_Assert(getDevice(oclinfo, CVCL_DEVICE_TYPE_ALL) > 0);
|
||||
oclinfo[0].impl->setDevice(0, 0, 0);
|
||||
clCxt.get()->impl = oclinfo[0].impl->copy();
|
||||
|
||||
*((volatile int*)&val) = 1;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,275 @@
|
||||
/*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
|
||||
// Sen Liu, swjtuls1987@126.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 oclMaterials 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*/
|
||||
|
||||
#ifndef WAVE_SIZE
|
||||
#define WAVE_SIZE 1
|
||||
#endif
|
||||
|
||||
int calc_lut(__local int* smem, int val, int tid)
|
||||
{
|
||||
smem[tid] = val;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid == 0)
|
||||
{
|
||||
for (int i = 1; i < 256; ++i)
|
||||
{
|
||||
smem[i] += smem[i - 1];
|
||||
}
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
return smem[tid];
|
||||
}
|
||||
|
||||
#ifdef CPU
|
||||
void reduce(volatile __local int* smem, int val, int tid)
|
||||
{
|
||||
smem[tid] = val;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 128)
|
||||
{
|
||||
smem[tid] = val += smem[tid + 128];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 64)
|
||||
{
|
||||
smem[tid] = val += smem[tid + 64];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 32)
|
||||
{
|
||||
smem[tid] += smem[tid + 32];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 16)
|
||||
{
|
||||
smem[tid] += smem[tid + 16];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 8)
|
||||
{
|
||||
smem[tid] += smem[tid + 8];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 4)
|
||||
{
|
||||
smem[tid] += smem[tid + 4];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 2)
|
||||
{
|
||||
smem[tid] += smem[tid + 2];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 1)
|
||||
{
|
||||
smem[256] = smem[tid] + smem[tid + 1];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
#else
|
||||
void reduce(__local volatile int* smem, int val, int tid)
|
||||
{
|
||||
smem[tid] = val;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 128)
|
||||
{
|
||||
smem[tid] = val += smem[tid + 128];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 64)
|
||||
{
|
||||
smem[tid] = val += smem[tid + 64];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (tid < 32)
|
||||
{
|
||||
smem[tid] += smem[tid + 32];
|
||||
#if WAVE_SIZE < 32
|
||||
} barrier(CLK_LOCAL_MEM_FENCE);
|
||||
if (tid < 16) {
|
||||
#endif
|
||||
smem[tid] += smem[tid + 16];
|
||||
#if WAVE_SIZE < 16
|
||||
} barrier(CLK_LOCAL_MEM_FENCE);
|
||||
if (tid < 8) {
|
||||
#endif
|
||||
smem[tid] += smem[tid + 8];
|
||||
smem[tid] += smem[tid + 4];
|
||||
smem[tid] += smem[tid + 2];
|
||||
smem[tid] += smem[tid + 1];
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
__kernel void calcLut(__global __const uchar * src, __global uchar * lut,
|
||||
const int srcStep, const int dstStep,
|
||||
const int2 tileSize, const int tilesX,
|
||||
const int clipLimit, const float lutScale)
|
||||
{
|
||||
__local int smem[512];
|
||||
|
||||
const int tx = get_group_id(0);
|
||||
const int ty = get_group_id(1);
|
||||
const unsigned int tid = get_local_id(1) * get_local_size(0)
|
||||
+ get_local_id(0);
|
||||
|
||||
smem[tid] = 0;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
for (int i = get_local_id(1); i < tileSize.y; i += get_local_size(1))
|
||||
{
|
||||
__global const uchar* srcPtr = src + mad24( ty * tileSize.y + i,
|
||||
srcStep, tx * tileSize.x );
|
||||
for (int j = get_local_id(0); j < tileSize.x; j += get_local_size(0))
|
||||
{
|
||||
const int data = srcPtr[j];
|
||||
atomic_inc(&smem[data]);
|
||||
}
|
||||
}
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
int tHistVal = smem[tid];
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (clipLimit > 0)
|
||||
{
|
||||
// clip histogram bar
|
||||
|
||||
int clipped = 0;
|
||||
if (tHistVal > clipLimit)
|
||||
{
|
||||
clipped = tHistVal - clipLimit;
|
||||
tHistVal = clipLimit;
|
||||
}
|
||||
|
||||
// find number of overall clipped samples
|
||||
|
||||
reduce(smem, clipped, tid);
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
#ifdef CPU
|
||||
clipped = smem[256];
|
||||
#else
|
||||
clipped = smem[0];
|
||||
#endif
|
||||
|
||||
// broadcast evaluated value
|
||||
|
||||
__local int totalClipped;
|
||||
|
||||
if (tid == 0)
|
||||
totalClipped = clipped;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
// redistribute clipped samples evenly
|
||||
|
||||
int redistBatch = totalClipped / 256;
|
||||
tHistVal += redistBatch;
|
||||
|
||||
int residual = totalClipped - redistBatch * 256;
|
||||
if (tid < residual)
|
||||
++tHistVal;
|
||||
}
|
||||
|
||||
const int lutVal = calc_lut(smem, tHistVal, tid);
|
||||
uint ires = (uint)convert_int_rte(lutScale * lutVal);
|
||||
lut[(ty * tilesX + tx) * dstStep + tid] =
|
||||
convert_uchar(clamp(ires, (uint)0, (uint)255));
|
||||
}
|
||||
|
||||
__kernel void transform(__global __const uchar * src,
|
||||
__global uchar * dst,
|
||||
__global uchar * lut,
|
||||
const int srcStep, const int dstStep, const int lutStep,
|
||||
const int cols, const int rows,
|
||||
const int2 tileSize,
|
||||
const int tilesX, const int tilesY)
|
||||
{
|
||||
const int x = get_global_id(0);
|
||||
const int y = get_global_id(1);
|
||||
|
||||
if (x >= cols || y >= rows)
|
||||
return;
|
||||
|
||||
const float tyf = (convert_float(y) / tileSize.y) - 0.5f;
|
||||
int ty1 = convert_int_rtn(tyf);
|
||||
int ty2 = ty1 + 1;
|
||||
const float ya = tyf - ty1;
|
||||
ty1 = max(ty1, 0);
|
||||
ty2 = min(ty2, tilesY - 1);
|
||||
|
||||
const float txf = (convert_float(x) / tileSize.x) - 0.5f;
|
||||
int tx1 = convert_int_rtn(txf);
|
||||
int tx2 = tx1 + 1;
|
||||
const float xa = txf - tx1;
|
||||
tx1 = max(tx1, 0);
|
||||
tx2 = min(tx2, tilesX - 1);
|
||||
|
||||
const int srcVal = src[mad24(y, srcStep, x)];
|
||||
|
||||
float res = 0;
|
||||
|
||||
res += lut[mad24(ty1 * tilesX + tx1, lutStep, srcVal)] * ((1.0f - xa) * (1.0f - ya));
|
||||
res += lut[mad24(ty1 * tilesX + tx2, lutStep, srcVal)] * ((xa) * (1.0f - ya));
|
||||
res += lut[mad24(ty2 * tilesX + tx1, lutStep, srcVal)] * ((1.0f - xa) * (ya));
|
||||
res += lut[mad24(ty2 * tilesX + tx2, lutStep, srcVal)] * ((xa) * (ya));
|
||||
|
||||
uint ires = (uint)convert_int_rte(res);
|
||||
dst[mad24(y, dstStep, x)] = convert_uchar(clamp(ires, (uint)0, (uint)255));
|
||||
}
|
||||
@@ -23,6 +23,7 @@
|
||||
// Rock Li, Rock.Li@amd.com
|
||||
// Wu Zailong, bullet@yeah.net
|
||||
// Xu Pang, pangxu010@163.com
|
||||
// Sen Liu, swjtuls1987@126.com
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
@@ -1393,6 +1394,46 @@ TEST_P(calcHist, Mat)
|
||||
EXPECT_MAT_NEAR(dst_hist, cpu_hist, 0.0);
|
||||
}
|
||||
}
|
||||
///////////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
// CLAHE
|
||||
namespace
|
||||
{
|
||||
IMPLEMENT_PARAM_CLASS(ClipLimit, double)
|
||||
}
|
||||
|
||||
PARAM_TEST_CASE(CLAHE, cv::Size, ClipLimit)
|
||||
{
|
||||
cv::Size size;
|
||||
double clipLimit;
|
||||
|
||||
cv::Mat src;
|
||||
cv::Mat dst_gold;
|
||||
|
||||
cv::ocl::oclMat g_src;
|
||||
cv::ocl::oclMat g_dst;
|
||||
|
||||
virtual void SetUp()
|
||||
{
|
||||
size = GET_PARAM(0);
|
||||
clipLimit = GET_PARAM(1);
|
||||
|
||||
cv::RNG &rng = TS::ptr()->get_rng();
|
||||
src = randomMat(rng, size, CV_8UC1, 0, 256, false);
|
||||
g_src.upload(src);
|
||||
}
|
||||
};
|
||||
|
||||
TEST_P(CLAHE, Accuracy)
|
||||
{
|
||||
cv::Ptr<cv::ocl::CLAHE> clahe = cv::ocl::createCLAHE(clipLimit);
|
||||
clahe->apply(g_src, g_dst);
|
||||
cv::Mat dst(g_dst);
|
||||
|
||||
cv::Ptr<cv::CLAHE> clahe_gold = cv::createCLAHE(clipLimit);
|
||||
clahe_gold->apply(src, dst_gold);
|
||||
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
|
||||
}
|
||||
|
||||
///////////////////////////Convolve//////////////////////////////////
|
||||
PARAM_TEST_CASE(ConvolveTestBase, MatType, bool)
|
||||
@@ -1643,6 +1684,10 @@ INSTANTIATE_TEST_CASE_P(histTestBase, calcHist, Combine(
|
||||
ONE_TYPE(CV_32SC1) //no use
|
||||
));
|
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, CLAHE, Combine(
|
||||
Values(cv::Size(128, 128), cv::Size(113, 113), cv::Size(1300, 1300)),
|
||||
Values(0.0, 40.0)));
|
||||
|
||||
//INSTANTIATE_TEST_CASE_P(ConvolveTestBase, Convolve, Combine(
|
||||
// Values(CV_32FC1, CV_32FC1),
|
||||
// Values(false))); // Values(false) is the reserved parameter
|
||||
|
||||
@@ -64,10 +64,6 @@ struct DistIdxPair
|
||||
|
||||
struct MatchPairsBody : ParallelLoopBody
|
||||
{
|
||||
MatchPairsBody(const MatchPairsBody& other)
|
||||
: matcher(other.matcher), features(other.features),
|
||||
pairwise_matches(other.pairwise_matches), near_pairs(other.near_pairs) {}
|
||||
|
||||
MatchPairsBody(FeaturesMatcher &_matcher, const std::vector<ImageFeatures> &_features,
|
||||
std::vector<MatchesInfo> &_pairwise_matches, std::vector<std::pair<int,int> > &_near_pairs)
|
||||
: matcher(_matcher), features(_features),
|
||||
|
||||
@@ -69,13 +69,13 @@ struct CalcRotation
|
||||
K_from(0,0) = cameras[edge.from].focal;
|
||||
K_from(1,1) = cameras[edge.from].focal * cameras[edge.from].aspect;
|
||||
K_from(0,2) = cameras[edge.from].ppx;
|
||||
K_from(0,2) = cameras[edge.from].ppy;
|
||||
K_from(1,2) = cameras[edge.from].ppy;
|
||||
|
||||
Mat_<double> K_to = Mat::eye(3, 3, CV_64F);
|
||||
K_to(0,0) = cameras[edge.to].focal;
|
||||
K_to(1,1) = cameras[edge.to].focal * cameras[edge.to].aspect;
|
||||
K_to(0,2) = cameras[edge.to].ppx;
|
||||
K_to(0,2) = cameras[edge.to].ppy;
|
||||
K_to(1,2) = cameras[edge.to].ppy;
|
||||
|
||||
Mat R = K_from.inv() * pairwise_matches[pair_idx].H.inv() * K_to;
|
||||
cameras[edge.to].R = cameras[edge.from].R * R;
|
||||
|
||||
@@ -165,7 +165,8 @@ PERF_TEST_P(Path_Idx_Cn_NPoints_WSize_Deriv, OpticalFlowPyrLK_self, testing::Com
|
||||
declare.in(pyramid1, pyramid2, inPoints).out(outPoints);
|
||||
declare.time(400);
|
||||
|
||||
TEST_CYCLE()
|
||||
int runs = 3;
|
||||
TEST_CYCLE_MULTIRUN(runs)
|
||||
{
|
||||
calcOpticalFlowPyrLK(pyramid1, pyramid2, inPoints, outPoints, status, err,
|
||||
Size(winSize, winSize), maxLevel, criteria,
|
||||
@@ -217,4 +218,4 @@ PERF_TEST_P(Path_Win_Deriv_Border_Reuse, OpticalFlowPyrLK_pyr, testing::Combine(
|
||||
}
|
||||
|
||||
SANITY_CHECK(pyramid);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -360,6 +360,9 @@ Mat estimateGlobalMotionRansac(
|
||||
const int npoints = points0.getMat().checkVector(2);
|
||||
CV_Assert(points1.getMat().checkVector(2) == npoints);
|
||||
|
||||
if (npoints < params.size)
|
||||
return Mat::eye(3, 3, CV_32F);
|
||||
|
||||
const Point2f *points0_ = points0.getMat().ptr<Point2f>();
|
||||
const Point2f *points1_ = points1.getMat().ptr<Point2f>();
|
||||
const int niters = params.niters();
|
||||
@@ -678,6 +681,8 @@ Mat KeypointBasedMotionEstimator::estimate(const Mat &frame0, const Mat &frame1,
|
||||
{
|
||||
// find keypoints
|
||||
detector_->detect(frame0, keypointsPrev_);
|
||||
if (keypointsPrev_.empty())
|
||||
return Mat::eye(3, 3, CV_32F);
|
||||
|
||||
// extract points from keypoints
|
||||
pointsPrev_.resize(keypointsPrev_.size());
|
||||
|
||||
Reference in New Issue
Block a user