generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Merge pull request #1914 from SpecLad:merge-2.4

Browse Source
This commit is contained in:
Roman Donchenko
2013-12-04 16:30:55 +04:00
committed by OpenCV Buildbot
parent f9190cc662 d58cd9851f
commit 0736ede7e5
125 changed files with 3283 additions and 304 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
modules/androidcamera/camera_wrapper/camera_wrapper.cpp
View File

@@ -686,7 +686,7 @@ void CameraHandler::closeCameraConnect()
camera->stopPreview();
#if defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3) || defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) \
|| defined(ANDROID_r4_3_0) || defined(ANDROID_r4_3_0)
|| defined(ANDROID_r4_3_0) || defined(ANDROID_r4_4_0)
camera->setPreviewCallbackFlags(CAMERA_FRAME_CALLBACK_FLAG_NOOP);
#endif
camera->disconnect();

2
modules/contrib/doc/facerec/src/facerec_eigenfaces.cpp
View File

@@ -70,7 +70,7 @@ int main(int argc, const char *argv[]) {
cout << "usage: " << argv[0] << " <csv.ext> <output_folder> " << endl;
exit(1);
}
string output_folder;
string output_folder = ".";
if (argc == 3) {
output_folder = string(argv[2]);
}

2
modules/contrib/doc/facerec/src/facerec_fisherfaces.cpp
View File

@@ -70,7 +70,7 @@ int main(int argc, const char *argv[]) {
cout << "usage: " << argv[0] << " <csv.ext> <output_folder> " << endl;
exit(1);
}
string output_folder;
string output_folder = ".";
if (argc == 3) {
output_folder = string(argv[2]);
}

2
modules/contrib/doc/facerec/src/facerec_save_load.cpp
View File

@@ -70,7 +70,7 @@ int main(int argc, const char *argv[]) {
cout << "usage: " << argv[0] << " <csv.ext> <output_folder> " << endl;
exit(1);
}
string output_folder;
string output_folder = ".";
if (argc == 3) {
output_folder = string(argv[2]);
}

4
modules/contrib/src/facerec.cpp
View File

@@ -602,8 +602,8 @@ inline void elbp_(InputArray _src, OutputArray _dst, int radius, int neighbors)
dst.setTo(0);
for(int n=0; n<neighbors; n++) {
// sample points
float x = static_cast<float>(-radius * sin(2.0*CV_PI*n/static_cast<float>(neighbors)));
float y = static_cast<float>(radius * cos(2.0*CV_PI*n/static_cast<float>(neighbors)));
float x = static_cast<float>(radius * cos(2.0*CV_PI*n/static_cast<float>(neighbors)));
float y = static_cast<float>(-radius * sin(2.0*CV_PI*n/static_cast<float>(neighbors)));
// relative indices
int fx = static_cast<int>(floor(x));
int fy = static_cast<int>(floor(y));

2
modules/core/doc/basic_structures.rst
View File

@@ -3122,5 +3122,5 @@ The above methods are usually enough for users. If you want to make your own alg
* Make a class and specify ``Algorithm`` as its base class.
* The algorithm parameters should be the class members. See ``Algorithm::get()`` for the list of possible types of the parameters.
* Add public virtual method ``AlgorithmInfo* info() const;`` to your class.
* Add constructor function, ``AlgorithmInfo`` instance and implement the ``info()`` method. The simplest way is to take http://code.opencv.org/projects/opencv/repository/revisions/master/entry/modules/ml/src/ml_init.cpp as the reference and modify it according to the list of your parameters.
* Add constructor function, ``AlgorithmInfo`` instance and implement the ``info()`` method. The simplest way is to take https://github.com/Itseez/opencv/tree/master/modules/ml/src/ml_init.cpp as the reference and modify it according to the list of your parameters.
* Add some public function (e.g. ``initModule_<mymodule>()``) that calls info() of your algorithm and put it into the same source file as ``info()`` implementation. This is to force C++ linker to include this object file into the target application. See ``Algorithm::create()`` for details.

13
modules/core/src/command_line_parser.cpp
View File

@@ -211,16 +211,15 @@ CommandLineParser::CommandLineParser(int argc, const char* const argv[], const S
}
impl->apply_params(k_v[0], k_v[1]);
}
else if (s.length() > 2 && s[0] == '-' && s[1] == '-')
{
impl->apply_params(s.substr(2), "true");
}
else if (s.length() > 1 && s[0] == '-')
{
for (int h = 0; h < 2; h++)
{
if (s[0] == '-')
s = s.substr(1, s.length() - 1);
}
impl->apply_params(s, "true");
impl->apply_params(s.substr(1), "true");
}
else if (s[0] != '-')
else
{
impl->apply_params(jj, s);
jj++;

2
modules/cuda/CMakeLists.txt
View File

@@ -1,4 +1,4 @@
if(ANDROID OR IOS)
if(IOS)
ocv_module_disable(cuda)
endif()

2
modules/cuda/doc/introduction.rst
View File

@@ -58,4 +58,4 @@ While developing algorithms for multiple GPUs, note a data passing overhead. For
3. Merge the results into a single disparity map.
With this algorithm, a dual GPU gave a 180% performance increase comparing to the single Fermi GPU. For a source code example, see http://code.opencv.org/projects/opencv/repository/revisions/master/entry/samples/gpu/.
With this algorithm, a dual GPU gave a 180% performance increase comparing to the single Fermi GPU. For a source code example, see https://github.com/Itseez/opencv/tree/master/samples/gpu/.

2
modules/cudaarithm/CMakeLists.txt
View File

@@ -1,4 +1,4 @@
if(ANDROID OR IOS)
if(IOS)
ocv_module_disable(cudaarithm)
endif()

2
modules/cudabgsegm/CMakeLists.txt
View File

@@ -1,4 +1,4 @@
if(ANDROID OR IOS)
if(IOS)
ocv_module_disable(cudabgsegm)
endif()

314
modules/cudabgsegm/perf/perf_bgsegm.cpp
View File

@@ -89,6 +89,8 @@ DEF_PARAM_TEST_1(Video, string);
PERF_TEST_P(Video, FGDStatModel,
Values(string("gpu/video/768x576.avi")))
{
const int numIters = 10;
declare.time(60);
const string inputFile = perf::TestBase::getDataPath(GetParam());
@@ -107,18 +109,36 @@ PERF_TEST_P(Video, FGDStatModel,
cv::Ptr<cv::cuda::BackgroundSubtractorFGD> d_fgd = cv::cuda::createBackgroundSubtractorFGD();
d_fgd->apply(d_frame, foreground);
for (int i = 0; i < 10; ++i)
int i = 0;
// collect performance data
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
d_frame.upload(frame);
startTimer(); next();
startTimer();
if(!next())
break;
d_fgd->apply(d_frame, foreground);
stopTimer();
}
// process last frame in sequence to get data for sanity test
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
d_frame.upload(frame);
d_fgd->apply(d_frame, foreground);
}
CUDA_SANITY_CHECK(foreground, 1e-2, ERROR_RELATIVE);
#ifdef HAVE_OPENCV_CUDAIMGPROC
@@ -134,18 +154,36 @@ PERF_TEST_P(Video, FGDStatModel,
IplImage ipl_frame = frame;
cv::Ptr<CvBGStatModel> model(cvCreateFGDStatModel(&ipl_frame));
for (int i = 0; i < 10; ++i)
int i = 0;
// collect performance data
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
ipl_frame = frame;
startTimer(); next();
startTimer();
if(!next())
break;
cvUpdateBGStatModel(&ipl_frame, model);
stopTimer();
}
// process last frame in sequence to get data for sanity test
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
ipl_frame = frame;
cvUpdateBGStatModel(&ipl_frame, model);
}
const cv::Mat background = cv::cvarrToMat(model->background);
const cv::Mat foreground = cv::cvarrToMat(model->foreground);
@@ -171,6 +209,8 @@ PERF_TEST_P(Video_Cn_LearningRate, MOG,
CUDA_CHANNELS_1_3_4,
Values(0.0, 0.01)))
{
const int numIters = 10;
const string inputFile = perf::TestBase::getDataPath(GET_PARAM(0));
const int cn = GET_PARAM(1);
const float learningRate = static_cast<float>(GET_PARAM(2));
@@ -202,7 +242,10 @@ PERF_TEST_P(Video_Cn_LearningRate, MOG,
d_mog->apply(d_frame, foreground, learningRate);
for (int i = 0; i < 10; ++i)
int i = 0;
// collect performance data
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
@@ -219,21 +262,17 @@ PERF_TEST_P(Video_Cn_LearningRate, MOG,
d_frame.upload(frame);
startTimer(); next();
startTimer();
if(!next())
break;
d_mog->apply(d_frame, foreground, learningRate);
stopTimer();
}
CUDA_SANITY_CHECK(foreground);
}
else
{
cv::Ptr<cv::BackgroundSubtractor> mog = cv::createBackgroundSubtractorMOG();
cv::Mat foreground;
mog->apply(frame, foreground, learningRate);
for (int i = 0; i < 10; ++i)
// process last frame in sequence to get data for sanity test
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
@@ -248,11 +287,66 @@ PERF_TEST_P(Video_Cn_LearningRate, MOG,
cv::swap(temp, frame);
}
startTimer(); next();
d_frame.upload(frame);
d_mog->apply(d_frame, foreground, learningRate);
}
CUDA_SANITY_CHECK(foreground);
}
else
{
cv::Ptr<cv::BackgroundSubtractor> mog = cv::createBackgroundSubtractorMOG();
cv::Mat foreground;
mog->apply(frame, foreground, learningRate);
int i = 0;
// collect performance data
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
if (cn != 3)
{
cv::Mat temp;
if (cn == 1)
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);
else
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA);
cv::swap(temp, frame);
}
startTimer();
if(!next())
break;
mog->apply(frame, foreground, learningRate);
stopTimer();
}
// process last frame in sequence to get data for sanity test
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
if (cn != 3)
{
cv::Mat temp;
if (cn == 1)
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);
else
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA);
cv::swap(temp, frame);
}
mog->apply(frame, foreground, learningRate);
}
CPU_SANITY_CHECK(foreground);
}
}
@@ -266,10 +360,12 @@ PERF_TEST_P(Video_Cn_LearningRate, MOG,
DEF_PARAM_TEST(Video_Cn, string, int);
PERF_TEST_P(Video_Cn, MOG2,
PERF_TEST_P(Video_Cn, DISABLED_MOG2,
Combine(Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"),
CUDA_CHANNELS_1_3_4))
{
const int numIters = 10;
const string inputFile = perf::TestBase::getDataPath(GET_PARAM(0));
const int cn = GET_PARAM(1);
@@ -301,7 +397,10 @@ PERF_TEST_P(Video_Cn, MOG2,
d_mog2->apply(d_frame, foreground);
for (int i = 0; i < 10; ++i)
int i = 0;
// collect performance data
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
@@ -318,23 +417,17 @@ PERF_TEST_P(Video_Cn, MOG2,
d_frame.upload(frame);
startTimer(); next();
startTimer();
if(!next())
break;
d_mog2->apply(d_frame, foreground);
stopTimer();
}
CUDA_SANITY_CHECK(foreground);
}
else
{
cv::Ptr<cv::BackgroundSubtractorMOG2> mog2 = cv::createBackgroundSubtractorMOG2();
mog2->setDetectShadows(false);
cv::Mat foreground;
mog2->apply(frame, foreground);
for (int i = 0; i < 10; ++i)
// process last frame in sequence to get data for sanity test
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
@@ -349,11 +442,68 @@ PERF_TEST_P(Video_Cn, MOG2,
cv::swap(temp, frame);
}
startTimer(); next();
d_frame.upload(frame);
d_mog2->apply(d_frame, foreground);
}
CUDA_SANITY_CHECK(foreground);
}
else
{
cv::Ptr<cv::BackgroundSubtractorMOG2> mog2 = cv::createBackgroundSubtractorMOG2();
mog2->setDetectShadows(false);
cv::Mat foreground;
mog2->apply(frame, foreground);
int i = 0;
// collect performance data
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
if (cn != 3)
{
cv::Mat temp;
if (cn == 1)
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);
else
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA);
cv::swap(temp, frame);
}
startTimer();
if(!next())
break;
mog2->apply(frame, foreground);
stopTimer();
}
// process last frame in sequence to get data for sanity test
for (; i < numIters; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
if (cn != 3)
{
cv::Mat temp;
if (cn == 1)
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);
else
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA);
cv::swap(temp, frame);
}
mog2->apply(frame, foreground);
}
CPU_SANITY_CHECK(foreground);
}
}
@@ -455,6 +605,8 @@ PERF_TEST_P(Video_Cn_MaxFeatures, GMG,
CUDA_CHANNELS_1_3_4,
Values(20, 40, 60)))
{
const int numIters = 150;
const std::string inputFile = perf::TestBase::getDataPath(GET_PARAM(0));
const int cn = GET_PARAM(1);
const int maxFeatures = GET_PARAM(2);
@@ -486,7 +638,10 @@ PERF_TEST_P(Video_Cn_MaxFeatures, GMG,
d_gmg->apply(d_frame, foreground);
for (int i = 0; i < 150; ++i)
int i = 0;
// collect performance data
for (; i < numIters; ++i)
{
cap >> frame;
if (frame.empty())
@@ -508,24 +663,17 @@ PERF_TEST_P(Video_Cn_MaxFeatures, GMG,
d_frame.upload(frame);
startTimer(); next();
startTimer();
if(!next())
break;
d_gmg->apply(d_frame, foreground);
stopTimer();
}
CUDA_SANITY_CHECK(foreground);
}
else
{
cv::Mat foreground;
cv::Mat zeros(frame.size(), CV_8UC1, cv::Scalar::all(0));
cv::Ptr<cv::BackgroundSubtractorGMG> gmg = cv::createBackgroundSubtractorGMG();
gmg->setMaxFeatures(maxFeatures);
gmg->apply(frame, foreground);
for (int i = 0; i < 150; ++i)
// process last frame in sequence to get data for sanity test
for (; i < numIters; ++i)
{
cap >> frame;
if (frame.empty())
@@ -545,11 +693,79 @@ PERF_TEST_P(Video_Cn_MaxFeatures, GMG,
cv::swap(temp, frame);
}
startTimer(); next();
d_frame.upload(frame);
d_gmg->apply(d_frame, foreground);
}
CUDA_SANITY_CHECK(foreground);
}
else
{
cv::Mat foreground;
cv::Mat zeros(frame.size(), CV_8UC1, cv::Scalar::all(0));
cv::Ptr<cv::BackgroundSubtractorGMG> gmg = cv::createBackgroundSubtractorGMG();
gmg->setMaxFeatures(maxFeatures);
gmg->apply(frame, foreground);
int i = 0;
// collect performance data
for (; i < numIters; ++i)
{
cap >> frame;
if (frame.empty())
{
cap.release();
cap.open(inputFile);
cap >> frame;
}
if (cn != 3)
{
cv::Mat temp;
if (cn == 1)
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);
else
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA);
cv::swap(temp, frame);
}
startTimer();
if(!next())
break;
gmg->apply(frame, foreground);
stopTimer();
}
// process last frame in sequence to get data for sanity test
for (; i < numIters; ++i)
{
cap >> frame;
if (frame.empty())
{
cap.release();
cap.open(inputFile);
cap >> frame;
}
if (cn != 3)
{
cv::Mat temp;
if (cn == 1)
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);
else
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA);
cv::swap(temp, frame);
}
gmg->apply(frame, foreground);
}
CPU_SANITY_CHECK(foreground);
}
}

2
modules/cudacodec/CMakeLists.txt
View File

@@ -1,4 +1,4 @@
if(ANDROID OR IOS OR APPLE)
if(IOS OR APPLE)
ocv_module_disable(cudacodec)
endif()

2
modules/cudafeatures2d/CMakeLists.txt
View File

@@ -1,4 +1,4 @@
if(ANDROID OR IOS)
if(IOS)
ocv_module_disable(cudafeatures2d)
endif()

2
modules/cudafilters/CMakeLists.txt
View File

@@ -1,4 +1,4 @@
if(ANDROID OR IOS)
if(IOS)
ocv_module_disable(cudafilters)
endif()

2
modules/cudaimgproc/CMakeLists.txt
View File

@@ -1,4 +1,4 @@
if(ANDROID OR IOS)
if(IOS)
ocv_module_disable(cudaimgproc)
endif()

2
modules/cudaoptflow/CMakeLists.txt
View File

@@ -1,4 +1,4 @@
if(ANDROID OR IOS)
if(IOS)
ocv_module_disable(cudaoptflow)
endif()

2
modules/cudastereo/CMakeLists.txt
View File

@@ -1,4 +1,4 @@
if(ANDROID OR IOS)
if(IOS)
ocv_module_disable(cudastereo)
endif()

2
modules/cudawarping/CMakeLists.txt
View File

@@ -1,4 +1,4 @@
if(ANDROID OR IOS)
if(IOS)
ocv_module_disable(cudawarping)
endif()

6
modules/highgui/doc/reading_and_writing_images_and_video.rst
View File

@@ -250,7 +250,7 @@ VideoCapture constructors.
.. ocv:cfunction:: CvCapture* cvCaptureFromCAM( int device )
.. ocv:cfunction:: CvCapture* cvCaptureFromFile( const char* filename )
:param filename: name of the opened video file (eg. video.avi) or image sequence (eg. img%02d.jpg)
:param filename: name of the opened video file (eg. video.avi) or image sequence (eg. img_%02d.jpg, which will read samples like img_00.jpg, img_01.jpg, img_02.jpg, ...)
:param device: id of the opened video capturing device (i.e. a camera index). If there is a single camera connected, just pass 0.
@@ -267,7 +267,7 @@ Open video file or a capturing device for video capturing
.. ocv:pyfunction:: cv2.VideoCapture.open(filename) -> retval
.. ocv:pyfunction:: cv2.VideoCapture.open(device) -> retval
:param filename: name of the opened video file (eg. video.avi) or image sequence (eg. img%02d.jpg)
:param filename: name of the opened video file (eg. video.avi) or image sequence (eg. img_%02d.jpg, which will read samples like img_00.jpg, img_01.jpg, img_02.jpg, ...)
:param device: id of the opened video capturing device (i.e. a camera index).
@@ -313,7 +313,7 @@ The methods/functions grab the next frame from video file or camera and return t
The primary use of the function is in multi-camera environments, especially when the cameras do not have hardware synchronization. That is, you call ``VideoCapture::grab()`` for each camera and after that call the slower method ``VideoCapture::retrieve()`` to decode and get frame from each camera. This way the overhead on demosaicing or motion jpeg decompression etc. is eliminated and the retrieved frames from different cameras will be closer in time.
Also, when a connected camera is multi-head (for example, a stereo camera or a Kinect device), the correct way of retrieving data from it is to call `VideoCapture::grab` first and then call :ocv:func:`VideoCapture::retrieve` one or more times with different values of the ``channel`` parameter. See http://code.opencv.org/projects/opencv/repository/revisions/master/entry/samples/cpp/kinect_maps.cpp
Also, when a connected camera is multi-head (for example, a stereo camera or a Kinect device), the correct way of retrieving data from it is to call `VideoCapture::grab` first and then call :ocv:func:`VideoCapture::retrieve` one or more times with different values of the ``channel`` parameter. See https://github.com/Itseez/opencv/tree/master/samples/cpp/openni_capture.cpp
VideoCapture::retrieve

2
modules/highgui/doc/user_interface.rst
View File

@@ -203,7 +203,7 @@ Sets mouse handler for the specified window
:param winname: Window name
:param onMouse: Mouse callback. See OpenCV samples, such as http://code.opencv.org/projects/opencv/repository/revisions/master/entry/samples/cpp/ffilldemo.cpp, on how to specify and use the callback.
:param onMouse: Mouse callback. See OpenCV samples, such as https://github.com/Itseez/opencv/tree/master/samples/cpp/ffilldemo.cpp, on how to specify and use the callback.
:param userdata: The optional parameter passed to the callback.

2
modules/imgproc/doc/structural_analysis_and_shape_descriptors.rst
View File

@@ -236,7 +236,7 @@ Approximates a polygonal curve(s) with the specified precision.
The functions ``approxPolyDP`` approximate a curve or a polygon with another curve/polygon with less vertices so that the distance between them is less or equal to the specified precision. It uses the Douglas-Peucker algorithm
http://en.wikipedia.org/wiki/Ramer-Douglas-Peucker_algorithm
See http://code.opencv.org/projects/opencv/repository/revisions/master/entry/samples/cpp/contours.cpp for the function usage model.
See https://github.com/Itseez/opencv/tree/master/samples/cpp/contours2.cpp for the function usage model.
ApproxChains

2
modules/objdetect/doc/cascade_classification.rst
View File

@@ -21,7 +21,7 @@ The word "cascade" in the classifier name means that the resultant classifier co
The feature used in a particular classifier is specified by its shape (1a, 2b etc.), position within the region of interest and the scale (this scale is not the same as the scale used at the detection stage, though these two scales are multiplied). For example, in the case of the third line feature (2c) the response is calculated as the difference between the sum of image pixels under the rectangle covering the whole feature (including the two white stripes and the black stripe in the middle) and the sum of the image pixels under the black stripe multiplied by 3 in order to compensate for the differences in the size of areas. The sums of pixel values over a rectangular regions are calculated rapidly using integral images (see below and the :ocv:func:`integral` description).
To see the object detector at work, have a look at the facedetect demo:
http://code.opencv.org/projects/opencv/repository/revisions/master/entry/samples/cpp/dbt_face_detection.cpp
https://github.com/Itseez/opencv/tree/master/samples/cpp/dbt_face_detection.cpp
The following reference is for the detection part only. There is a separate application called ``opencv_traincascade`` that can train a cascade of boosted classifiers from a set of samples.

2
modules/ocl/include/opencv2/ocl/matrix_operations.hpp
View File

@@ -44,6 +44,8 @@
#ifndef __OPENCV_OCL_MATRIX_OPERATIONS_HPP__
#define __OPENCV_OCL_MATRIX_OPERATIONS_HPP__
#include "opencv2/ocl.hpp"
namespace cv
{

45
modules/ocl/src/cl_context.cpp
View File

@@ -189,11 +189,8 @@ static bool parseOpenCLDeviceConfiguration(const std::string& configurationStr,
return true;
}
static bool __deviceSelected = false;
static bool selectOpenCLDevice()
{
__deviceSelected = true;
std::string platform;
std::vector<std::string> deviceTypes;
std::string deviceName;
@@ -528,26 +525,38 @@ private:
static ContextImpl* currentContext = NULL;
static bool __deviceSelected = false;
Context* Context::getContext()
{
if (currentContext == NULL)
{
if (!__initialized || !__deviceSelected)
static bool defaultInitiaization = false;
if (!defaultInitiaization)
{
cv::AutoLock lock(getInitializationMutex());
if (!__initialized)
try
{
if (initializeOpenCLDevices() == 0)
if (!__initialized)
{
CV_Error(Error::OpenCLInitError, "OpenCL not available");
if (initializeOpenCLDevices() == 0)
{
CV_Error(Error::OpenCLInitError, "OpenCL not available");
}
}
if (!__deviceSelected)
{
if (!selectOpenCLDevice())
{
CV_Error(Error::OpenCLInitError, "Can't select OpenCL device");
}
}
defaultInitiaization = true;
}
if (!__deviceSelected)
catch (...)
{
if (!selectOpenCLDevice())
{
CV_Error(Error::OpenCLInitError, "Can't select OpenCL device");
}
defaultInitiaization = true;
throw;
}
}
CV_Assert(currentContext != NULL);
@@ -744,10 +753,16 @@ int getOpenCLDevices(std::vector<const DeviceInfo*> &devices, int deviceType, co
void setDevice(const DeviceInfo* info)
{
if (!__deviceSelected)
try
{
ContextImpl::setContext(info);
__deviceSelected = true;
ContextImpl::setContext(info);
}
catch (...)
{
__deviceSelected = true;
throw;
}
}
bool supportsFeature(FEATURE_TYPE featureType)

15
modules/ocl/src/cl_operations.cpp
View File

@@ -192,6 +192,7 @@ void openCLMallocPitchEx(Context *ctx, void **dev_ptr, size_t *pitch,
clFinish(getClCommandQueue(ctx));
#endif
CheckBuffers data(mainBuffer, size, widthInBytes, height);
cv::AutoLock lock(getInitializationMutex());
__check_buffers.insert(std::pair<cl_mem, CheckBuffers>((cl_mem)*dev_ptr, data));
}
#endif
@@ -253,10 +254,17 @@ void openCLFree(void *devPtr)
bool failBefore = false, failAfter = false;
#endif
CheckBuffers data;
std::map<cl_mem, CheckBuffers>::iterator i = __check_buffers.find((cl_mem)devPtr);
if (i != __check_buffers.end())
{
data = i->second;
cv::AutoLock lock(getInitializationMutex());
std::map<cl_mem, CheckBuffers>::iterator i = __check_buffers.find((cl_mem)devPtr);
if (i != __check_buffers.end())
{
data = i->second;
__check_buffers.erase(i);
}
}
if (data.mainBuffer != NULL)
{
#ifdef CHECK_MEMORY_CORRUPTION
Context* ctx = Context::getContext();
std::vector<uchar> checkBefore(__memory_corruption_guard_bytes);
@@ -286,7 +294,6 @@ void openCLFree(void *devPtr)
clFinish(getClCommandQueue(ctx));
#endif
openCLSafeCall(clReleaseMemObject(data.mainBuffer));
__check_buffers.erase(i);
}
#if defined(CHECK_MEMORY_CORRUPTION)
if (failBefore)

2
modules/ocl/src/haar.cpp
View File

@@ -923,7 +923,7 @@ void OclCascadeClassifier::detectMultiScale(oclMat &gimg, CV_OUT std::vector<cv:
//use known local data stride to precalulate indexes
int DATA_SIZE_X = (localThreads[0]+cascade->orig_window_size.width);
// check that maximal value is less than maximal unsigned short
assert(DATA_SIZE_X*cascade->orig_window_size.height+cascade->orig_window_size.width < USHRT_MAX);
assert(DATA_SIZE_X*cascade->orig_window_size.height+cascade->orig_window_size.width < (int)USHRT_MAX);
for(int i = 0;i<nodenum;++i)
{//process each node from classifier
struct NodePK

4
modules/ocl/src/mssegmentation.cpp
View File

@@ -42,6 +42,10 @@
//
//M*/
#if (__GNUC__ == 4) && (__GNUC_MINOR__ == 6)
# pragma GCC diagnostic ignored "-Warray-bounds"
#endif
#include "precomp.hpp"
#include "opencl_kernels.hpp"

33
modules/ocl/src/opencl/imgproc_calcHarris.cl
View File

@@ -47,6 +47,13 @@
/////////////////////////////////Macro for border type////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////////////////
#if defined (DOUBLE_SUPPORT) && defined (cl_khr_fp64)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#define FPTYPE double
#else
#define FPTYPE float
#endif
#ifdef BORDER_CONSTANT
#elif defined BORDER_REPLICATE
#define EXTRAPOLATE(x, maxV) \
@@ -116,7 +123,7 @@ __kernel void calcHarris(__global const float *Dx, __global const float *Dy, __g
int dst_startY = (gY << 1) + dst_y_off;
float dx_data[ksY+1],dy_data[ksY+1], data[3][ksY+1];
__local float temp[6][THREADS];
__local FPTYPE temp[6][THREADS];
#ifdef BORDER_CONSTANT
for (int i=0; i < ksY+1; i++)
@@ -136,7 +143,7 @@ __kernel void calcHarris(__global const float *Dx, __global const float *Dy, __g
data[2][i] = dy_data[i] * dy_data[i];
}
#else
int clamped_col = min(dst_cols, col);
int clamped_col = min(2*dst_cols, col);
for (int i=0; i < ksY+1; i++)
{
int dx_selected_row = dx_startY+i, dx_selected_col = dx_startX+clamped_col;
@@ -154,7 +161,7 @@ __kernel void calcHarris(__global const float *Dx, __global const float *Dy, __g
data[2][i] = dy_data[i] * dy_data[i];
}
#endif
float sum0 = 0.0f, sum1 = 0.0f, sum2 = 0.0f;
FPTYPE sum0 = 0.0f, sum1 = 0.0f, sum2 = 0.0f;
for (int i=1; i < ksY; i++)
{
sum0 += data[0][i];
@@ -162,16 +169,16 @@ __kernel void calcHarris(__global const float *Dx, __global const float *Dy, __g
sum2 += data[2][i];
}
float sum01 = sum0 + data[0][0];
float sum02 = sum0 + data[0][ksY];
FPTYPE sum01 = sum0 + data[0][0];
FPTYPE sum02 = sum0 + data[0][ksY];
temp[0][col] = sum01;
temp[1][col] = sum02;
float sum11 = sum1 + data[1][0];
float sum12 = sum1 + data[1][ksY];
FPTYPE sum11 = sum1 + data[1][0];
FPTYPE sum12 = sum1 + data[1][ksY];
temp[2][col] = sum11;
temp[3][col] = sum12;
float sum21 = sum2 + data[2][0];
float sum22 = sum2 + data[2][ksY];
FPTYPE sum21 = sum2 + data[2][0];
FPTYPE sum22 = sum2 + data[2][ksY];
temp[4][col] = sum21;
temp[5][col] = sum22;
barrier(CLK_LOCAL_MEM_FENCE);
@@ -184,8 +191,14 @@ __kernel void calcHarris(__global const float *Dx, __global const float *Dy, __g
int till = (ksX + 1)%2;
float tmp_sum[6] = { 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f };
for (int k=0; k<6; k++)
{
FPTYPE temp_sum = 0;
for (int i=-anX; i<=anX - till; i++)
tmp_sum[k] += temp[k][col+i];
{
temp_sum += temp[k][col+i];
}
tmp_sum[k] = temp_sum;
}
if (posX < dst_cols && (posY) < dst_rows)
{

24
modules/ocl/test/test_imgproc.cpp
View File

@@ -254,8 +254,28 @@ OCL_TEST_P(CornerMinEigenVal, Mat)
}
////////////////////////////////cornerHarris//////////////////////////////////////////
struct CornerHarris :
public ImgprocTestBase
{
void Near(double threshold = 0.0)
{
Mat whole, roi;
gdst_whole.download(whole);
gdst_roi.download(roi);
typedef CornerTestBase CornerHarris;
absdiff(whole, dst_whole, whole);
absdiff(roi, dst_roi, roi);
divide(whole, dst_whole, whole);
divide(roi, dst_roi, roi);
absdiff(dst_whole, dst_whole, dst_whole);
absdiff(dst_roi, dst_roi, dst_roi);
EXPECT_MAT_NEAR(dst_whole, whole, threshold);
EXPECT_MAT_NEAR(dst_roi, roi, threshold);
}
};
OCL_TEST_P(CornerHarris, Mat)
{
@@ -269,7 +289,7 @@ OCL_TEST_P(CornerHarris, Mat)
cornerHarris(src_roi, dst_roi, blockSize, apertureSize, k, borderType);
ocl::cornerHarris(gsrc_roi, gdst_roi, blockSize, apertureSize, k, borderType);
Near(1e-5, true);
Near(1e-5);
}
}