1
0
mirror of https://github.com/opencv/opencv.git synced 2026-07-29 23:33:05 +04:00

Deleted all trailing whitespace.

This commit is contained in:
Roman Donchenko
2013-08-21 16:44:09 +04:00
parent 0d8cb2e319
commit f55740da70
193 changed files with 1685 additions and 1685 deletions
@@ -1481,7 +1481,7 @@ Reconstructs points by triangulation.
:param points4D: 4xN array of reconstructed points in homogeneous coordinates.
The function reconstructs 3-dimensional points (in homogeneous coordinates) by using their observations with a stereo camera. Projections matrices can be obtained from :ocv:func:`stereoRectify`.
The function reconstructs 3-dimensional points (in homogeneous coordinates) by using their observations with a stereo camera. Projections matrices can be obtained from :ocv:func:`stereoRectify`.
.. seealso::
@@ -4,19 +4,19 @@ Changelog
Release 0.05
------------
This library is now included in the official OpenCV distribution (from 2.4 on).
This library is now included in the official OpenCV distribution (from 2.4 on).
The :ocv:class`FaceRecognizer` is now an :ocv:class:`Algorithm`, which better fits into the overall
OpenCV API.
OpenCV API.
To reduce the confusion on user side and minimize my work, libfacerec and OpenCV
have been synchronized and are now based on the same interfaces and implementation.
To reduce the confusion on user side and minimize my work, libfacerec and OpenCV
have been synchronized and are now based on the same interfaces and implementation.
The library now has an extensive documentation:
* The API is explained in detail and with a lot of code examples.
* The face recognition guide I had written for Python and GNU Octave/MATLAB has been adapted to the new OpenCV C++ ``cv::FaceRecognizer``.
* The face recognition guide I had written for Python and GNU Octave/MATLAB has been adapted to the new OpenCV C++ ``cv::FaceRecognizer``.
* A tutorial for gender classification with Fisherfaces.
* A tutorial for face recognition in videos (e.g. webcam).
* A tutorial for face recognition in videos (e.g. webcam).
Release highlights
@@ -27,8 +27,8 @@ Release highlights
Release 0.04
------------
This version is fully Windows-compatible and works with OpenCV 2.3.1. Several
bugfixes, but none influenced the recognition rate.
This version is fully Windows-compatible and works with OpenCV 2.3.1. Several
bugfixes, but none influenced the recognition rate.
Release highlights
++++++++++++++++++
@@ -40,9 +40,9 @@ Release highlights
Release 0.03
------------
Reworked the library to provide separate implementations in cpp files, because
it's the preferred way of contributing OpenCV libraries. This means the library
is not header-only anymore. Slight API changes were done, please see the
Reworked the library to provide separate implementations in cpp files, because
it's the preferred way of contributing OpenCV libraries. This means the library
is not header-only anymore. Slight API changes were done, please see the
documentation for details.
Release highlights
@@ -55,9 +55,9 @@ Release highlights
Release 0.02
------------
Reworked the library to provide separate implementations in cpp files, because
it's the preferred way of contributing OpenCV libraries. This means the library
is not header-only anymore. Slight API changes were done, please see the
Reworked the library to provide separate implementations in cpp files, because
it's the preferred way of contributing OpenCV libraries. This means the library
is not header-only anymore. Slight API changes were done, please see the
documentation for details.
Release highlights
@@ -80,7 +80,7 @@ Release highlights
* Eigenfaces [TP91]_
* Fisherfaces [BHK97]_
* Local Binary Patterns Histograms [AHP04]_
* Added persistence facilities to store the models with a common API.
* Unit Tests (using `gtest <http://code.google.com/p/googletest/>`_).
* Providing a CMakeLists.txt to enable easy cross-platform building.
@@ -201,7 +201,7 @@ For the first source code example, I'll go through it with you. I am first givin
.. literalinclude:: src/facerec_eigenfaces.cpp
:language: cpp
:linenos:
The source code for this demo application is also available in the ``src`` folder coming with this documentation:
* :download:`src/facerec_eigenfaces.cpp <src/facerec_eigenfaces.cpp>`
@@ -6,7 +6,7 @@ Introduction
Saving and loading a :ocv:class:`FaceRecognizer` is very important. Training a FaceRecognizer can be a very time-intense task, plus it's often impossible to ship the whole face database to the user of your product. The task of saving and loading a FaceRecognizer is easy with :ocv:class:`FaceRecognizer`. You only have to call :ocv:func:`FaceRecognizer::load` for loading and :ocv:func:`FaceRecognizer::save` for saving a :ocv:class:`FaceRecognizer`.
I'll adapt the Eigenfaces example from the :doc:`../facerec_tutorial`: Imagine we want to learn the Eigenfaces of the `AT&T Facedatabase <http://www.cl.cam.ac.uk/research/dtg/attarchive/facedatabase.html>`_, store the model to a YAML file and then load it again.
I'll adapt the Eigenfaces example from the :doc:`../facerec_tutorial`: Imagine we want to learn the Eigenfaces of the `AT&T Facedatabase <http://www.cl.cam.ac.uk/research/dtg/attarchive/facedatabase.html>`_, store the model to a YAML file and then load it again.
From the loaded model, we'll get a prediction, show the mean, Eigenfaces and the image reconstruction.
@@ -111,7 +111,7 @@ An example. If the haar-cascade is at ``C:/opencv/data/haarcascades/haarcascade_
facerec_video.exe C:/opencv/data/haarcascades/haarcascade_frontalface_default.xml C:/facerec/data/celebrities.txt 1
That's it.
That's it.
Results
-------
+1 -1
View File
@@ -306,7 +306,7 @@ void FaceRecognizer::update(InputArrayOfArrays src, InputArray labels ) {
dynamic_cast<LBPH*>(this)->update( src, labels );
return;
}
string error_msg = format("This FaceRecognizer (%s) does not support updating, you have to use FaceRecognizer::train to update it.", this->name().c_str());
CV_Error(CV_StsNotImplemented, error_msg);
}
+2 -2
View File
@@ -553,7 +553,7 @@ Range
-----
.. ocv:class:: Range
Template class specifying a continuous subsequence (slice) of a sequence.
Template class specifying a continuous subsequence (slice) of a sequence.
::
@@ -773,7 +773,7 @@ Mat
---
.. ocv:class:: Mat
OpenCV C++ n-dimensional dense array class
OpenCV C++ n-dimensional dense array class
::
class CV_EXPORTS Mat
+2 -2
View File
@@ -80,8 +80,8 @@ Splits an element set into equivalency classes.
:param vec: Set of elements stored as a vector.
:param labels: Output vector of labels. It contains as many elements as ``vec``. Each label ``labels[i]`` is a 0-based cluster index of ``vec[i]`` .
:param labels: Output vector of labels. It contains as many elements as ``vec``. Each label ``labels[i]`` is a 0-based cluster index of ``vec[i]`` .
:param predicate: Equivalence predicate (pointer to a boolean function of two arguments or an instance of the class that has the method ``bool operator()(const _Tp& a, const _Tp& b)`` ). The predicate returns ``true`` when the elements are certainly in the same class, and returns ``false`` if they may or may not be in the same class.
The generic function ``partition`` implements an
+2 -2
View File
@@ -416,8 +416,8 @@ The number of pixels along the line is stored in ``LineIterator::count`` . The m
for(int i = 0; i < it.count; i++, ++it)
buf[i] = *(const Vec3b)*it;
// alternative way of iterating through the line
// alternative way of iterating through the line
for(int i = 0; i < it2.count; i++, ++it2)
{
Vec3b val = img.at<Vec3b>(it2.pos());
+2 -2
View File
@@ -91,8 +91,8 @@ you can use::
Ptr<T> ptr = new T(...);
That is, ``Ptr<T> ptr`` encapsulates a pointer to a ``T`` instance and a reference counter associated with the pointer. See the
:ocv:class:`Ptr`
That is, ``Ptr<T> ptr`` encapsulates a pointer to a ``T`` instance and a reference counter associated with the pointer. See the
:ocv:class:`Ptr`
description for details.
.. _AutomaticAllocation:
@@ -3002,55 +3002,55 @@ static inline void read(const FileNode& node, string& value, const string& defau
}
template<typename _Tp> static inline void read(const FileNode& node, Point_<_Tp>& value, const Point_<_Tp>& default_value)
{
{
vector<_Tp> temp; FileNodeIterator it = node.begin(); it >> temp;
value = temp.size() != 2 ? default_value : Point_<_Tp>(saturate_cast<_Tp>(temp[0]), saturate_cast<_Tp>(temp[1]));
}
template<typename _Tp> static inline void read(const FileNode& node, Point3_<_Tp>& value, const Point3_<_Tp>& default_value)
{
{
vector<_Tp> temp; FileNodeIterator it = node.begin(); it >> temp;
value = temp.size() != 3 ? default_value : Point3_<_Tp>(saturate_cast<_Tp>(temp[0]), saturate_cast<_Tp>(temp[1]),
saturate_cast<_Tp>(temp[2]));
}
template<typename _Tp> static inline void read(const FileNode& node, Size_<_Tp>& value, const Size_<_Tp>& default_value)
{
{
vector<_Tp> temp; FileNodeIterator it = node.begin(); it >> temp;
value = temp.size() != 2 ? default_value : Size_<_Tp>(saturate_cast<_Tp>(temp[0]), saturate_cast<_Tp>(temp[1]));
}
template<typename _Tp> static inline void read(const FileNode& node, Complex<_Tp>& value, const Complex<_Tp>& default_value)
{
{
vector<_Tp> temp; FileNodeIterator it = node.begin(); it >> temp;
value = temp.size() != 2 ? default_value : Complex<_Tp>(saturate_cast<_Tp>(temp[0]), saturate_cast<_Tp>(temp[1]));
}
template<typename _Tp> static inline void read(const FileNode& node, Rect_<_Tp>& value, const Rect_<_Tp>& default_value)
{
{
vector<_Tp> temp; FileNodeIterator it = node.begin(); it >> temp;
value = temp.size() != 4 ? default_value : Rect_<_Tp>(saturate_cast<_Tp>(temp[0]), saturate_cast<_Tp>(temp[1]),
value = temp.size() != 4 ? default_value : Rect_<_Tp>(saturate_cast<_Tp>(temp[0]), saturate_cast<_Tp>(temp[1]),
saturate_cast<_Tp>(temp[2]), saturate_cast<_Tp>(temp[3]));
}
template<typename _Tp, int cn> static inline void read(const FileNode& node, Vec<_Tp, cn>& value, const Vec<_Tp, cn>& default_value)
{
{
vector<_Tp> temp; FileNodeIterator it = node.begin(); it >> temp;
value = temp.size() != cn ? default_value : Vec<_Tp, cn>(&temp[0]);
}
template<typename _Tp> static inline void read(const FileNode& node, Scalar_<_Tp>& value, const Scalar_<_Tp>& default_value)
{
{
vector<_Tp> temp; FileNodeIterator it = node.begin(); it >> temp;
value = temp.size() != 4 ? default_value : Scalar_<_Tp>(saturate_cast<_Tp>(temp[0]), saturate_cast<_Tp>(temp[1]),
saturate_cast<_Tp>(temp[2]), saturate_cast<_Tp>(temp[3]));
}
static inline void read(const FileNode& node, Range& value, const Range& default_value)
{
Point2i temp(value.start, value.end); const Point2i default_temp = Point2i(default_value.start, default_value.end);
{
Point2i temp(value.start, value.end); const Point2i default_temp = Point2i(default_value.start, default_value.end);
read(node, temp, default_temp);
value.start = temp.x; value.end = temp.y;
value.start = temp.x; value.end = temp.y;
}
CV_EXPORTS_W void read(const FileNode& node, Mat& mat, const Mat& default_mat=Mat() );
+4 -4
View File
@@ -1252,14 +1252,14 @@ static void arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
Mat src1 = _src1.getMat(), src2 = _src2.getMat();
bool haveMask = !_mask.empty();
bool reallocate = false;
bool src1Scalar = checkScalar(src1, src2.type(), kind1, kind2);
bool src2Scalar = checkScalar(src2, src1.type(), kind2, kind1);
bool src1Scalar = checkScalar(src1, src2.type(), kind1, kind2);
bool src2Scalar = checkScalar(src2, src1.type(), kind2, kind1);
if( (kind1 == kind2 || src1.channels() == 1) && src1.dims <= 2 && src2.dims <= 2 &&
src1.size() == src2.size() && src1.type() == src2.type() &&
!haveMask && ((!_dst.fixedType() && (dtype < 0 || CV_MAT_DEPTH(dtype) == src1.depth())) ||
(_dst.fixedType() && _dst.type() == _src1.type())) &&
(_dst.fixedType() && _dst.type() == _src1.type())) &&
((src1Scalar && src2Scalar) || (!src1Scalar && !src2Scalar)) )
{
_dst.create(src1.size(), src1.type());
+9 -9
View File
@@ -453,7 +453,7 @@ cv::Scalar cv::sum( InputArray _src )
{
Mat src = _src.getMat();
int k, cn = src.channels(), depth = src.depth();
#if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
size_t total_size = src.total();
int rows = src.size[0], cols = (int)(total_size/rows);
@@ -462,7 +462,7 @@ cv::Scalar cv::sum( InputArray _src )
IppiSize sz = { cols, rows };
int type = src.type();
typedef IppStatus (CV_STDCALL* ippiSumFunc)(const void*, int, IppiSize, double *, int);
ippiSumFunc ippFunc =
ippiSumFunc ippFunc =
type == CV_8UC1 ? (ippiSumFunc)ippiSum_8u_C1R :
type == CV_8UC3 ? (ippiSumFunc)ippiSum_8u_C3R :
type == CV_8UC4 ? (ippiSumFunc)ippiSum_8u_C4R :
@@ -490,8 +490,8 @@ cv::Scalar cv::sum( InputArray _src )
}
}
}
#endif
#endif
SumFunc func = getSumFunc(depth);
CV_Assert( cn <= 4 && func != 0 );
@@ -565,7 +565,7 @@ cv::Scalar cv::mean( InputArray _src, InputArray _mask )
CV_Assert( mask.empty() || mask.type() == CV_8U );
int k, cn = src.channels(), depth = src.depth();
#if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
size_t total_size = src.total();
int rows = src.size[0], cols = (int)(total_size/rows);
@@ -576,7 +576,7 @@ cv::Scalar cv::mean( InputArray _src, InputArray _mask )
if( !mask.empty() )
{
typedef IppStatus (CV_STDCALL* ippiMaskMeanFuncC1)(const void *, int, void *, int, IppiSize, Ipp64f *);
ippiMaskMeanFuncC1 ippFuncC1 =
ippiMaskMeanFuncC1 ippFuncC1 =
type == CV_8UC1 ? (ippiMaskMeanFuncC1)ippiMean_8u_C1MR :
type == CV_16UC1 ? (ippiMaskMeanFuncC1)ippiMean_16u_C1MR :
type == CV_32FC1 ? (ippiMaskMeanFuncC1)ippiMean_32f_C1MR :
@@ -590,7 +590,7 @@ cv::Scalar cv::mean( InputArray _src, InputArray _mask )
}
}
typedef IppStatus (CV_STDCALL* ippiMaskMeanFuncC3)(const void *, int, void *, int, IppiSize, int, Ipp64f *);
ippiMaskMeanFuncC3 ippFuncC3 =
ippiMaskMeanFuncC3 ippFuncC3 =
type == CV_8UC3 ? (ippiMaskMeanFuncC3)ippiMean_8u_C3CMR :
type == CV_16UC3 ? (ippiMaskMeanFuncC3)ippiMean_16u_C3CMR :
type == CV_32FC3 ? (ippiMaskMeanFuncC3)ippiMean_32f_C3CMR :
@@ -609,7 +609,7 @@ cv::Scalar cv::mean( InputArray _src, InputArray _mask )
else
{
typedef IppStatus (CV_STDCALL* ippiMeanFunc)(const void*, int, IppiSize, double *, int);
ippiMeanFunc ippFunc =
ippiMeanFunc ippFunc =
type == CV_8UC1 ? (ippiMeanFunc)ippiMean_8u_C1R :
type == CV_8UC3 ? (ippiMeanFunc)ippiMean_8u_C3R :
type == CV_8UC4 ? (ippiMeanFunc)ippiMean_8u_C4R :
@@ -639,7 +639,7 @@ cv::Scalar cv::mean( InputArray _src, InputArray _mask )
}
}
#endif
SumFunc func = getSumFunc(depth);
CV_Assert( cn <= 4 && func != 0 );
+1 -1
View File
@@ -405,7 +405,7 @@ protected:
Vec<int, 5> v1(15, 16, 17, 18, 19), ov1;
Scalar sc1(20.0, 21.1, 22.2, 23.3), osc1;
Range g1(7, 8), og1;
FileStorage fs(fname, FileStorage::WRITE);
fs << "mi" << mi;
fs << "mv" << mv;
+1 -1
View File
@@ -2457,7 +2457,7 @@ TEST(Core_Invert, small)
{
cv::Mat a = (cv::Mat_<float>(3,3) << 2.42104644730331, 1.81444796521479, -3.98072565304758, 0, 7.08389214348967e-3, 5.55326770986007e-3, 0,0, 7.44556154284261e-3);
//cv::randu(a, -1, 1);
cv::Mat b = a.t()*a;
cv::Mat c, i = Mat_<float>::eye(3, 3);
cv::invert(b, c, cv::DECOMP_LU); //std::cout << b*c << std::endl;
@@ -40,7 +40,7 @@ Lixin Fan, Jutta Willamowski, Cedric Bray, 2004. ::
BOWTrainer::add
-------------------
Adds descriptors to a training set.
Adds descriptors to a training set.
.. ocv:function:: void BOWTrainer::add( const Mat& descriptors )
@@ -66,7 +66,7 @@ Returns the count of all descriptors stored in the training set.
BOWTrainer::cluster
-----------------------
Clusters train descriptors.
Clusters train descriptors.
.. ocv:function:: Mat BOWTrainer::cluster() const
@@ -116,7 +116,7 @@ Class to compute an image descriptor using the *bag of visual words*. Such a com
#. Compute descriptors for a given image and its keypoints set.
#. Find the nearest visual words from the vocabulary for each keypoint descriptor.
#. Compute the bag-of-words image descriptor as is a normalized histogram of vocabulary words encountered in the image. The ``i``-th bin of the histogram is a frequency of ``i``-th word of the vocabulary in the given image.
The class declaration is the following: ::
class BOWImgDescriptorExtractor
+2 -2
View File
@@ -178,10 +178,10 @@ void BOWImgDescriptorExtractor::compute( const Mat& image, vector<KeyPoint>& key
// Normalize image descriptor.
imgDescriptor /= descriptors.rows;
// Add the descriptors of image keypoints
if (_descriptors) {
*_descriptors = descriptors.clone();
*_descriptors = descriptors.clone();
}
}
+1 -1
View File
@@ -258,7 +258,7 @@ struct IntersectAreaCounter
{
CV_Assert( miny < maxy );
CV_Assert( dr > FLT_EPSILON );
int temp_bua = bua, temp_bna = bna;
for( int i = range.begin(); i != range.end(); i++ )
{
@@ -68,11 +68,11 @@ The method constructs a fast search structure from a set of features using the s
* **branching** The branching factor to use for the hierarchical k-means tree
* **iterations** The maximum number of iterations to use in the k-means clustering stage when building the k-means tree. A value of -1 used here means that the k-means clustering should be iterated until convergence
* **iterations** The maximum number of iterations to use in the k-means clustering stage when building the k-means tree. A value of -1 used here means that the k-means clustering should be iterated until convergence
* **centers_init** The algorithm to use for selecting the initial centers when performing a k-means clustering step. The possible values are ``CENTERS_RANDOM`` (picks the initial cluster centers randomly), ``CENTERS_GONZALES`` (picks the initial centers using Gonzales' algorithm) and ``CENTERS_KMEANSPP`` (picks the initial centers using the algorithm suggested in arthur_kmeanspp_2007 )
* **centers_init** The algorithm to use for selecting the initial centers when performing a k-means clustering step. The possible values are ``CENTERS_RANDOM`` (picks the initial cluster centers randomly), ``CENTERS_GONZALES`` (picks the initial centers using Gonzales' algorithm) and ``CENTERS_KMEANSPP`` (picks the initial centers using the algorithm suggested in arthur_kmeanspp_2007 )
* **cb_index** This parameter (cluster boundary index) influences the way exploration is performed in the hierarchical kmeans tree. When ``cb_index`` is zero the next kmeans domain to be explored is chosen to be the one with the closest center. A value greater then zero also takes into account the size of the domain.
* **cb_index** This parameter (cluster boundary index) influences the way exploration is performed in the hierarchical kmeans tree. When ``cb_index`` is zero the next kmeans domain to be explored is chosen to be the one with the closest center. A value greater then zero also takes into account the size of the domain.
*
**CompositeIndexParams** When using a parameters object of this type the index created combines the randomized kd-trees and the hierarchical k-means tree. ::
@@ -122,16 +122,16 @@ The method constructs a fast search structure from a set of features using the s
..
* **target_precision** Is a number between 0 and 1 specifying the percentage of the approximate nearest-neighbor searches that return the exact nearest-neighbor. Using a higher value for this parameter gives more accurate results, but the search takes longer. The optimum value usually depends on the application.
* **target_precision** Is a number between 0 and 1 specifying the percentage of the approximate nearest-neighbor searches that return the exact nearest-neighbor. Using a higher value for this parameter gives more accurate results, but the search takes longer. The optimum value usually depends on the application.
* **build_weight** Specifies the importance of the index build time raported to the nearest-neighbor search time. In some applications it's acceptable for the index build step to take a long time if the subsequent searches in the index can be performed very fast. In other applications it's required that the index be build as fast as possible even if that leads to slightly longer search times.
* **build_weight** Specifies the importance of the index build time raported to the nearest-neighbor search time. In some applications it's acceptable for the index build step to take a long time if the subsequent searches in the index can be performed very fast. In other applications it's required that the index be build as fast as possible even if that leads to slightly longer search times.
* **memory_weight** Is used to specify the tradeoff between time (index build time and search time) and memory used by the index. A value less than 1 gives more importance to the time spent and a value greater than 1 gives more importance to the memory usage.
* **sample_fraction** Is a number between 0 and 1 indicating what fraction of the dataset to use in the automatic parameter configuration algorithm. Running the algorithm on the full dataset gives the most accurate results, but for very large datasets can take longer than desired. In such case using just a fraction of the data helps speeding up this algorithm while still giving good approximations of the optimum parameters.
* **sample_fraction** Is a number between 0 and 1 indicating what fraction of the dataset to use in the automatic parameter configuration algorithm. Running the algorithm on the full dataset gives the most accurate results, but for very large datasets can take longer than desired. In such case using just a fraction of the data helps speeding up this algorithm while still giving good approximations of the optimum parameters.
*
**SavedIndexParams** This object type is used for loading a previously saved index from the disk. ::
+1 -1
View File
@@ -43,7 +43,7 @@ if(HAVE_CUDA)
ocv_cuda_compile(cuda_objs ${lib_cuda} ${ncv_cuda})
set(cuda_link_libs ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
if(HAVE_CUFFT)
set(cuda_link_libs ${cuda_link_libs} ${CUDA_cufft_LIBRARY})
endif()
@@ -49,7 +49,7 @@ This means that the input left image is low textured.
* A basic stereo matching example can be found at opencv_source_code/samples/gpu/stereo_match.cpp
* A stereo matching example using several GPU's can be found at opencv_source_code/samples/gpu/stereo_multi.cpp
* A stereo matching example using several GPU's and driver API can be found at opencv_source_code/samples/gpu/driver_api_stereo_multi.cpp
gpu::StereoBM_GPU::StereoBM_GPU
-----------------------------------
Enables :ocv:class:`gpu::StereoBM_GPU` constructors.
+1 -1
View File
@@ -7,7 +7,7 @@ Video Analysis
* A general optical flow example can be found at opencv_source_code/samples/gpu/optical_flow.cpp
* A general optical flow example using the Nvidia API can be found at opencv_source_code/samples/gpu/opticalflow_nvidia_api.cpp
gpu::BroxOpticalFlow
--------------------
.. ocv:class:: gpu::BroxOpticalFlow
@@ -20,7 +20,7 @@ Reads an image from a buffer in memory.
:param buf: Input array or vector of bytes.
:param flags: The same flags as in :ocv:func:`imread` .
:param dst: The optional output placeholder for the decoded matrix. It can save the image reallocations when the function is called repeatedly for images of the same size.
The function reads an image from the specified buffer in the memory.
@@ -74,9 +74,9 @@ Loads an image from a file.
:param filename: Name of file to be loaded.
:param flags: Flags specifying the color type of a loaded image:
* CV_LOAD_IMAGE_ANYDEPTH - If set, return 16-bit/32-bit image when the input has the corresponding depth, otherwise convert it to 8-bit.
* CV_LOAD_IMAGE_COLOR - If set, always convert image to the color one
* CV_LOAD_IMAGE_GRAYSCALE - If set, always convert image to the grayscale one
+2 -2
View File
@@ -160,7 +160,7 @@ private:
};
class CvCapture_FFMPEG_proxy :
class CvCapture_FFMPEG_proxy :
public CvCapture
{
public:
@@ -224,7 +224,7 @@ CvCapture* cvCreateFileCapture_FFMPEG_proxy(const char * filename)
return 0;
}
class CvVideoWriter_FFMPEG_proxy :
class CvVideoWriter_FFMPEG_proxy :
public CvVideoWriter
{
public:
+2 -2
View File
@@ -345,7 +345,7 @@ class ImplMutex
public:
ImplMutex() { init(); }
~ImplMutex() { destroy(); }
void init();
void destroy();
@@ -450,7 +450,7 @@ void ImplMutex::init()
impl = (Impl*)malloc(sizeof(Impl));
impl->init();
}
void ImplMutex::destroy()
void ImplMutex::destroy()
{
impl->destroy();
free(impl);
+7 -7
View File
@@ -388,7 +388,7 @@ static CGFloat DegreesToRadians(CGFloat degrees) {return degrees * M_PI / 180;};
- (CVPixelBufferRef) pixelBufferFromCGImage: (CGImageRef) image
{
CGSize frameSize = CGSizeMake(CGImageGetWidth(image), CGImageGetHeight(image));
NSDictionary *options = [NSDictionary dictionaryWithObjectsAndKeys:
[NSNumber numberWithBool:NO], kCVPixelBufferCGImageCompatibilityKey,
@@ -399,23 +399,23 @@ static CGFloat DegreesToRadians(CGFloat degrees) {return degrees * M_PI / 180;};
frameSize.height, kCVPixelFormatType_32ARGB, (CFDictionaryRef) CFBridgingRetain(options),
&pxbuffer);
NSParameterAssert(status == kCVReturnSuccess && pxbuffer != NULL);
CVPixelBufferLockBaseAddress(pxbuffer, 0);
void *pxdata = CVPixelBufferGetBaseAddress(pxbuffer);
CGColorSpaceRef rgbColorSpace = CGColorSpaceCreateDeviceRGB();
CGContextRef context = CGBitmapContextCreate(pxdata, frameSize.width,
frameSize.height, 8, 4*frameSize.width, rgbColorSpace,
kCGImageAlphaPremultipliedFirst);
CGContextDrawImage(context, CGRectMake(0, 0, CGImageGetWidth(image),
CGImageGetHeight(image)), image);
CGColorSpaceRelease(rgbColorSpace);
CGContextRelease(context);
CVPixelBufferUnlockBaseAddress(pxbuffer, 0);
return pxbuffer;
}
+1 -1
View File
@@ -14,7 +14,7 @@ It has been tested with the motempl sample program
First Patch: August 24, 2004 Travis Wood TravisOCV@tkwood.com
For Release: OpenCV-Linux Beta4 opencv-0.9.6
Tested On: LMLBT44 with 8 video inputs
Problems? Post your questions at answers.opencv.org,
Problems? Post your questions at answers.opencv.org,
Report bugs at code.opencv.org,
Submit your fixes at https://github.com/Itseez/opencv/
Patched Comments:
+2 -2
View File
@@ -3026,7 +3026,7 @@ double CvCaptureFile_MSMF::getProperty(int property_id)
return ((double)captureFormats[captureFormatIndex].MF_MT_FRAME_RATE_NUMERATOR) /
((double)captureFormats[captureFormatIndex].MF_MT_FRAME_RATE_DENOMINATOR);
}
return -1;
}
@@ -3062,7 +3062,7 @@ IplImage* CvCaptureFile_MSMF::retrieveFrame(int)
if(RIOut && size == RIOut->getSize())
{
videoInput::processPixels(RIOut->getpPixels(), (unsigned char*)frame->imageData, width,
videoInput::processPixels(RIOut->getpPixels(), (unsigned char*)frame->imageData, width,
height, bytes, false, verticalFlip);
}
+4 -4
View File
@@ -413,9 +413,9 @@ int CvCaptureCAM::startCaptureDevice(int cameraNum) {
void CvCaptureCAM::setWidthHeight() {
NSAutoreleasePool* localpool = [[NSAutoreleasePool alloc] init];
[mCaptureSession stopRunning];
NSDictionary* pixelBufferOptions = [NSDictionary dictionaryWithObjectsAndKeys:
[NSNumber numberWithDouble:1.0*width], (id)kCVPixelBufferWidthKey,
[NSNumber numberWithDouble:1.0*height], (id)kCVPixelBufferHeightKey,
@@ -424,9 +424,9 @@ void CvCaptureCAM::setWidthHeight() {
nil];
[mCaptureDecompressedVideoOutput setPixelBufferAttributes:pixelBufferOptions];
[mCaptureSession startRunning];
grabFrame(60);
[localpool drain];
}
+5 -5
View File
@@ -14,7 +14,7 @@ It has been tested with the motempl sample program
First Patch: August 24, 2004 Travis Wood TravisOCV@tkwood.com
For Release: OpenCV-Linux Beta4 opencv-0.9.6
Tested On: LMLBT44 with 8 video inputs
Problems? Post your questions at answers.opencv.org,
Problems? Post your questions at answers.opencv.org,
Report bugs at code.opencv.org,
Submit your fixes at https://github.com/Itseez/opencv/
Patched Comments:
@@ -157,7 +157,7 @@ the symptoms were damaged image and 'Corrupt JPEG data: premature end of data se
prevents bad images in the first place
11th patch: April 2, 2013, Forrest Reiling forrest.reiling@gmail.com
Added v4l2 support for getting capture property CV_CAP_PROP_POS_MSEC.
Added v4l2 support for getting capture property CV_CAP_PROP_POS_MSEC.
Returns the millisecond timestamp of the last frame grabbed or 0 if no frames have been grabbed
Used to successfully synchonize 2 Logitech C310 USB webcams to within 16 ms of one another
@@ -1233,8 +1233,8 @@ static int read_frame_v4l2(CvCaptureCAM_V4L* capture) {
if (-1 == ioctl (capture->deviceHandle, VIDIOC_QBUF, &buf))
perror ("VIDIOC_QBUF");
//set timestamp in capture struct to be timestamp of most recent frame
capture->timestamp = buf.timestamp;
//set timestamp in capture struct to be timestamp of most recent frame
capture->timestamp = buf.timestamp;
return 1;
}
@@ -2327,7 +2327,7 @@ static double icvGetPropertyCAM_V4L (CvCaptureCAM_V4L* capture,
if (capture->FirstCapture) {
return 0;
} else {
return 1000 * capture->timestamp.tv_sec + ((double) capture->timestamp.tv_usec) / 1000;
return 1000 * capture->timestamp.tv_sec + ((double) capture->timestamp.tv_usec) / 1000;
}
break;
case CV_CAP_PROP_BRIGHTNESS:
+9 -9
View File
@@ -138,7 +138,7 @@ void CvCaptureCAM_XIMEA::close()
{
if(frame)
cvReleaseImage(&frame);
if(hmv)
{
xiStopAcquisition(hmv);
@@ -176,11 +176,11 @@ IplImage* CvCaptureCAM_XIMEA::retrieveFrame(int)
{
// update cvImage after format has changed
resetCvImage();
// copy pixel data
switch( image.frm)
{
case XI_MONO8 :
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;
@@ -210,15 +210,15 @@ void CvCaptureCAM_XIMEA::resetCvImage()
{
case XI_MONO8 :
case XI_RAW8 : frame = cvCreateImage(cvSize( image.width, image.height), IPL_DEPTH_8U, 1); break;
case XI_MONO16 :
case XI_MONO16 :
case XI_RAW16 : frame = cvCreateImage(cvSize( image.width, image.height), IPL_DEPTH_16U, 1); break;
case XI_RGB24 :
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);
}
/**********************************************************************************/
@@ -338,9 +338,9 @@ int CvCaptureCAM_XIMEA::getBpp()
{
case XI_MONO8 :
case XI_RAW8 : return 1;
case XI_MONO16 :
case XI_MONO16 :
case XI_RAW16 : return 2;
case XI_RGB24 :
case XI_RGB24 :
case XI_RGB_PLANAR : return 3;
case XI_RGB32 : return 4;
default :
@@ -348,4 +348,4 @@ int CvCaptureCAM_XIMEA::getBpp()
}
}
/**********************************************************************************/
/**********************************************************************************/
@@ -16,4 +16,4 @@ The license does not permit the following uses:
You may not use, or allow anyone else to use the icons to create pornographic, libelous, obscene, or defamatory material.
All icon files are provided "as is". You agree not to hold IconEden.com liable for any damages that may occur due to use, or inability to use, icons or image data from IconEden.com.
All icon files are provided "as is". You agree not to hold IconEden.com liable for any damages that may occur due to use, or inability to use, icons or image data from IconEden.com.
+8 -8
View File
@@ -44,21 +44,21 @@
#include "precomp.hpp"
UIImage* MatToUIImage(const cv::Mat& image) {
NSData *data = [NSData dataWithBytes:image.data
length:image.elemSize()*image.total()];
CGColorSpaceRef colorSpace;
if (image.elemSize() == 1) {
colorSpace = CGColorSpaceCreateDeviceGray();
} else {
colorSpace = CGColorSpaceCreateDeviceRGB();
}
CGDataProviderRef provider =
CGDataProviderCreateWithCFData((__bridge CFDataRef)data);
// Creating CGImage from cv::Mat
CGImageRef imageRef = CGImageCreate(image.cols,
image.rows,
@@ -73,14 +73,14 @@ UIImage* MatToUIImage(const cv::Mat& image) {
false,
kCGRenderingIntentDefault
);
// Getting UIImage from CGImage
UIImage *finalImage = [UIImage imageWithCGImage:imageRef];
CGImageRelease(imageRef);
CGDataProviderRelease(provider);
CGColorSpaceRelease(colorSpace);
return finalImage;
}
+1 -1
View File
@@ -2474,7 +2474,7 @@ void DefaultViewPort::saveView()
if (!fileName.isEmpty()) //save the picture
{
QString extension = fileName.right(3);
// Create a new pixmap to render the viewport into
QPixmap viewportPixmap(viewport()->size());
viewport()->render(&viewportPixmap);
+1 -1
View File
@@ -179,7 +179,7 @@ Compares two histograms.
* **CV_COMP_INTERSECT** Intersection
* **CV_COMP_BHATTACHARYYA** Bhattacharyya distance
* **CV_COMP_HELLINGER** Synonym for ``CV_COMP_BHATTACHARYYA``
The functions ``compareHist`` compare two dense or two sparse histograms using the specified method:
@@ -309,7 +309,7 @@ enum
// alpha premultiplication
CV_RGBA2mRGBA = 125,
CV_mRGBA2RGBA = 126,
CV_RGB2YUV_I420 = 127,
CV_BGR2YUV_I420 = 128,
CV_RGB2YUV_IYUV = CV_RGB2YUV_I420,
+1 -1
View File
@@ -3896,7 +3896,7 @@ void cv::cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
CV_Error( CV_StsBadArg, "Unsupported image depth" );
}
}
break;
break;
default:
CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );
}
+21 -21
View File
@@ -1149,11 +1149,11 @@ static bool IPPMorphReplicate(int op, const Mat &src, Mat &dst, const Mat &kerne
}
//DEPRECATED. Allocates and initializes morphology state structure for erosion or dilation operation.
typedef IppStatus (CV_STDCALL* ippiMorphologyInitAllocFunc)(int, const void*, IppiSize, IppiPoint, IppiMorphState **);
ippiMorphologyInitAllocFunc ippInitAllocFunc =
type == CV_8UC1 ? (ippiMorphologyInitAllocFunc)ippiMorphologyInitAlloc_8u_C1R :
type == CV_8UC3 ? (ippiMorphologyInitAllocFunc)ippiMorphologyInitAlloc_8u_C3R :
type == CV_8UC4 ? (ippiMorphologyInitAllocFunc)ippiMorphologyInitAlloc_8u_C4R :
type == CV_32FC1 ? (ippiMorphologyInitAllocFunc)ippiMorphologyInitAlloc_32f_C1R :
ippiMorphologyInitAllocFunc ippInitAllocFunc =
type == CV_8UC1 ? (ippiMorphologyInitAllocFunc)ippiMorphologyInitAlloc_8u_C1R :
type == CV_8UC3 ? (ippiMorphologyInitAllocFunc)ippiMorphologyInitAlloc_8u_C3R :
type == CV_8UC4 ? (ippiMorphologyInitAllocFunc)ippiMorphologyInitAlloc_8u_C4R :
type == CV_32FC1 ? (ippiMorphologyInitAllocFunc)ippiMorphologyInitAlloc_32f_C1R :
type == CV_32FC3 ? (ippiMorphologyInitAllocFunc)ippiMorphologyInitAlloc_32f_C3R :
type == CV_32FC4 ? (ippiMorphologyInitAllocFunc)ippiMorphologyInitAlloc_32f_C4R :
0;
@@ -1163,25 +1163,25 @@ static bool IPPMorphReplicate(int op, const Mat &src, Mat &dst, const Mat &kerne
{
case MORPH_DILATE:
{
ippFunc =
type == CV_8UC1 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_8u_C1R :
type == CV_8UC3 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_8u_C3R :
type == CV_8UC4 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_8u_C4R :
type == CV_32FC1 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_32f_C1R :
type == CV_32FC3 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_32f_C3R :
type == CV_32FC4 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_32f_C4R :
ippFunc =
type == CV_8UC1 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_8u_C1R :
type == CV_8UC3 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_8u_C3R :
type == CV_8UC4 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_8u_C4R :
type == CV_32FC1 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_32f_C1R :
type == CV_32FC3 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_32f_C3R :
type == CV_32FC4 ? (ippiMorphologyBorderReplicateFunc)ippiDilateBorderReplicate_32f_C4R :
0;
break;
}
case MORPH_ERODE:
{
ippFunc =
type == CV_8UC1 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_8u_C1R :
type == CV_8UC3 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_8u_C3R :
type == CV_8UC4 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_8u_C4R :
type == CV_32FC1 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_32f_C1R :
type == CV_32FC3 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_32f_C3R :
type == CV_32FC4 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_32f_C4R :
ippFunc =
type == CV_8UC1 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_8u_C1R :
type == CV_8UC3 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_8u_C3R :
type == CV_8UC4 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_8u_C4R :
type == CV_32FC1 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_32f_C1R :
type == CV_32FC3 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_32f_C3R :
type == CV_32FC4 ? (ippiMorphologyBorderReplicateFunc)ippiErodeBorderReplicate_32f_C4R :
0;
break;
}
@@ -1207,8 +1207,8 @@ static bool IPPMorphOp(int op, InputArray _src, OutputArray _dst,
int borderType, const Scalar &borderValue)
{
Mat src = _src.getMat(), kernel = _kernel.getMat();
if( !( src.depth() == CV_8U || src.depth() == CV_32F ) || ( iterations > 1 ) ||
!( borderType == cv::BORDER_REPLICATE || (borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue()) )
if( !( src.depth() == CV_8U || src.depth() == CV_32F ) || ( iterations > 1 ) ||
!( borderType == cv::BORDER_REPLICATE || (borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue()) )
|| !( op == MORPH_DILATE || op == MORPH_ERODE) )
return false;
if( borderType == cv::BORDER_CONSTANT )
+1 -1
View File
@@ -349,7 +349,7 @@ public:
int r = rgb[0];
int g = rgb[1];
int b = rgb[2];
uchar y = saturate_cast<uchar>((int)( 0.257f*r + 0.504f*g + 0.098f*b + 0.5f) + 16);
uchar u = saturate_cast<uchar>((int)(-0.148f*r - 0.291f*g + 0.439f*b + 0.5f) + 128);
uchar v = saturate_cast<uchar>((int)( 0.439f*r - 0.368f*g - 0.071f*b + 0.5f) + 128);
+41 -41
View File
@@ -11,10 +11,10 @@ using namespace cv;
static void throwJavaException(JNIEnv *env, const std::exception *e, const char *method) {
std::string what = "unknown exception";
jclass je = 0;
if(e) {
std::string exception_type = "std::exception";
if(dynamic_cast<const cv::Exception*>(e)) {
exception_type = "cv::Exception";
je = env->FindClass("org/opencv/core/CvException");
@@ -22,16 +22,16 @@ static void throwJavaException(JNIEnv *env, const std::exception *e, const char
what = exception_type + ": " + e->what();
}
if(!je) je = env->FindClass("java/lang/Exception");
env->ThrowNew(je, what.c_str());
LOGE("%s caught %s", method, what.c_str());
(void)method; // avoid "unused" warning
}
extern "C" {
//
// MatXXX::MatXXX()
@@ -69,7 +69,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1Mat__III
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -95,7 +95,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1Mat__DDI
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -122,7 +122,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1Mat__IIIDDDD
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -149,7 +149,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1Mat__DDIDDDD
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -176,7 +176,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1Mat__JIIII
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -198,7 +198,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1Mat__JII
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -224,7 +224,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1adjustROI
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -293,7 +293,7 @@ JNIEXPORT jint JNICALL Java_org_opencv_core_Mat_n_1channels
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -319,7 +319,7 @@ JNIEXPORT jint JNICALL Java_org_opencv_core_Mat_n_1checkVector__JIIZ
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -341,7 +341,7 @@ JNIEXPORT jint JNICALL Java_org_opencv_core_Mat_n_1checkVector__JII
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -363,7 +363,7 @@ JNIEXPORT jint JNICALL Java_org_opencv_core_Mat_n_1checkVector__JI
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -391,7 +391,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1clone
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -418,7 +418,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1col
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -445,7 +445,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1colRange
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -454,7 +454,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1colRange
//
// int Mat::dims()
//
JNIEXPORT jint JNICALL Java_org_opencv_core_Mat_n_1dims
(JNIEnv* env, jclass, jlong self);
@@ -471,7 +471,7 @@ JNIEXPORT jint JNICALL Java_org_opencv_core_Mat_n_1dims
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -497,7 +497,7 @@ JNIEXPORT jint JNICALL Java_org_opencv_core_Mat_n_1cols
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -690,7 +690,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1cross
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -733,7 +733,7 @@ JNIEXPORT jint JNICALL Java_org_opencv_core_Mat_n_1depth
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -760,7 +760,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1diag__JI
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -787,7 +787,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1diag__J
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -814,7 +814,7 @@ JNIEXPORT jdouble JNICALL Java_org_opencv_core_Mat_n_1dot
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -840,7 +840,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1elemSize
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -866,7 +866,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1elemSize1
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -892,7 +892,7 @@ JNIEXPORT jboolean JNICALL Java_org_opencv_core_Mat_n_1empty
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -918,7 +918,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1eye__III
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -945,7 +945,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1eye__DDI
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -972,7 +972,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1inv__JI
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -994,7 +994,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1inv__J
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -1020,7 +1020,7 @@ JNIEXPORT jboolean JNICALL Java_org_opencv_core_Mat_n_1isContinuous
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -1046,7 +1046,7 @@ JNIEXPORT jboolean JNICALL Java_org_opencv_core_Mat_n_1isSubmatrix
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -1101,7 +1101,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1mul__JJD
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -1253,7 +1253,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1reshape__JII
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -1276,7 +1276,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1reshape__JI
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -1303,7 +1303,7 @@ JNIEXPORT jlong JNICALL Java_org_opencv_core_Mat_n_1row
} catch (...) {
throwJavaException(env, 0, method_name);
}
return 0;
}
@@ -1487,7 +1487,7 @@ JNIEXPORT jdoubleArray JNICALL Java_org_opencv_core_Mat_n_1size
LOGD("%s", method_name);
Mat* me = (Mat*) self; //TODO: check for NULL
Size _retval_ = me->size( );
jdoubleArray _da_retval_ = env->NewDoubleArray(2);
jdoubleArray _da_retval_ = env->NewDoubleArray(2);
jdouble _tmp_retval_[2] = {_retval_.width, _retval_.height};
env->SetDoubleArrayRegion(_da_retval_, 0, 2, _tmp_retval_);
return _da_retval_;
+15 -15
View File
@@ -67,7 +67,7 @@ The following loss functions are implemented for regression problems:
:math:`L(y,f(x)) = \left\{ \begin{array}{lr}
\delta\cdot\left(|y-f(x)|-\dfrac{\delta}{2}\right) & : |y-f(x)|>\delta\\
\dfrac{1}{2}\cdot(y-f(x))^2 & : |y-f(x)|\leq\delta \end{array} \right.`,
where :math:`\delta` is the :math:`\alpha`-quantile estimation of the
:math:`|y-f(x)|`. In the current implementation :math:`\alpha=0.2`.
@@ -129,9 +129,9 @@ CvGBTreesParams::CvGBTreesParams
:param weak_count: Count of boosting algorithm iterations. ``weak_count*K`` is the total
count of trees in the GBT model, where ``K`` is the output classes count
(equal to one in case of a regression).
:param shrinkage: Regularization parameter (see :ref:`Training GBT`).
:param subsample_portion: Portion of the whole training set used for each algorithm iteration.
Subset is generated randomly. For more information see
http://www.salfordsystems.com/doc/StochasticBoostingSS.pdf.
@@ -139,7 +139,7 @@ CvGBTreesParams::CvGBTreesParams
:param max_depth: Maximal depth of each decision tree in the ensemble (see :ocv:class:`CvDTree`).
:param use_surrogates: If ``true``, surrogate splits are built (see :ocv:class:`CvDTree`).
By default the following constructor is used:
.. code-block:: cpp
@@ -178,7 +178,7 @@ Trains a Gradient boosted tree model.
.. ocv:function:: bool CvGBTrees::train(CvMLData* data, CvGBTreesParams params=CvGBTreesParams(), bool update=false)
.. ocv:pyfunction:: cv2.GBTrees.train(trainData, tflag, responses[, varIdx[, sampleIdx[, varType[, missingDataMask[, params[, update]]]]]]) -> retval
The first train method follows the common template (see :ocv:func:`CvStatModel::train`).
Both ``tflag`` values (``CV_ROW_SAMPLE``, ``CV_COL_SAMPLE``) are supported.
``trainData`` must be of the ``CV_32F`` type. ``responses`` must be a matrix of type
@@ -188,7 +188,7 @@ list of indices (``CV_32S``) or a mask (``CV_8U`` or ``CV_8S``). ``update`` is
a dummy parameter.
The second form of :ocv:func:`CvGBTrees::train` function uses :ocv:class:`CvMLData` as a
data set container. ``update`` is still a dummy parameter.
data set container. ``update`` is still a dummy parameter.
All parameters specific to the GBT model are passed into the training function
as a :ocv:class:`CvGBTreesParams` structure.
@@ -207,42 +207,42 @@ Predicts a response for an input sample.
:param sample: Input feature vector that has the same format as every training set
element. If not all the variables were actually used during training,
``sample`` contains forged values at the appropriate places.
:param missing: Missing values mask, which is a dimensional matrix of the same size as
``sample`` having the ``CV_8U`` type. ``1`` corresponds to the missing value
in the same position in the ``sample`` vector. If there are no missing values
in the feature vector, an empty matrix can be passed instead of the missing mask.
:param weakResponses: Matrix used to obtain predictions of all the trees.
The matrix has :math:`K` rows,
where :math:`K` is the count of output classes (1 for the regression case).
The matrix has as many columns as the ``slice`` length.
:param slice: Parameter defining the part of the ensemble used for prediction.
If ``slice = Range::all()``, all trees are used. Use this parameter to
get predictions of the GBT models with different ensemble sizes learning
only one model.
:param k: Number of tree ensembles built in case of the classification problem
(see :ref:`Training GBT`). Use this
parameter to change the output to sum of the trees' predictions in the
``k``-th ensemble only. To get the total GBT model prediction, ``k`` value
must be -1. For regression problems, ``k`` is also equal to -1.
The method predicts the response corresponding to the given sample
(see :ref:`Predicting with GBT`).
The result is either the class label or the estimated function value. The
:ocv:func:`CvGBTrees::predict` method enables using the parallel version of the GBT model
prediction if the OpenCV is built with the TBB library. In this case, predictions
of single trees are computed in a parallel fashion.
of single trees are computed in a parallel fashion.
CvGBTrees::clear
----------------
Clears the model.
.. ocv:function:: void CvGBTrees::clear()
.. ocv:pyfunction:: cv2.GBTrees.clear() -> None
The function deletes the data set information and all the weak models and sets all internal
@@ -257,7 +257,7 @@ Calculates a training or testing error.
.. ocv:function:: float CvGBTrees::calc_error( CvMLData* _data, int type, std::vector<float> *resp = 0 )
:param _data: Data set.
:param type: Parameter defining the error that should be computed: train (``CV_TRAIN_ERROR``) or test
(``CV_TEST_ERROR``).
+1 -1
View File
@@ -45,7 +45,7 @@ Trains the model.
:param updateBase: Specifies whether the model is trained from scratch (``update_base=false``), or it is updated using the new training data (``update_base=true``). In the latter case, the parameter ``maxK`` must not be larger than the original value.
The method trains the K-Nearest model. It follows the conventions of the generic :ocv:func:`CvStatModel::train` approach with the following limitations:
The method trains the K-Nearest model. It follows the conventions of the generic :ocv:func:`CvStatModel::train` approach with the following limitations:
* Only ``CV_ROW_SAMPLE`` data layout is supported.
* Input variables are all ordered.
+31 -31
View File
@@ -9,7 +9,7 @@ CvMLData
--------
.. ocv:class:: CvMLData
Class for loading the data from a ``.csv`` file.
Class for loading the data from a ``.csv`` file.
::
class CV_EXPORTS CvMLData
@@ -27,42 +27,42 @@ Class for loading the data from a ``.csv`` file.
void set_response_idx( int idx );
int get_response_idx() const;
void set_train_test_split( const CvTrainTestSplit * spl);
const CvMat* get_train_sample_idx() const;
const CvMat* get_test_sample_idx() const;
void mix_train_and_test_idx();
const CvMat* get_var_idx();
void change_var_idx( int vi, bool state );
const CvMat* get_var_types();
void set_var_types( const char* str );
int get_var_type( int var_idx ) const;
void change_var_type( int var_idx, int type);
void set_delimiter( char ch );
char get_delimiter() const;
void set_miss_ch( char ch );
char get_miss_ch() const;
const std::map<std::string, int>& get_class_labels_map() const;
protected:
...
protected:
...
};
CvMLData::read_csv
------------------
Reads the data set from a ``.csv``-like ``filename`` file and stores all read values in a matrix.
Reads the data set from a ``.csv``-like ``filename`` file and stores all read values in a matrix.
.. ocv:function:: int CvMLData::read_csv(const char* filename)
:param filename: The input file name
While reading the data, the method tries to define the type of variables (predictors and responses): ordered or categorical. If a value of the variable is not numerical (except for the label for a missing value), the type of the variable is set to ``CV_VAR_CATEGORICAL``. If all existing values of the variable are numerical, the type of the variable is set to ``CV_VAR_ORDERED``. So, the default definition of variables types works correctly for all cases except the case of a categorical variable with numerical class labels. In this case, the type ``CV_VAR_ORDERED`` is set. You should change the type to ``CV_VAR_CATEGORICAL`` using the method :ocv:func:`CvMLData::change_var_type`. For categorical variables, a common map is built to convert a string class label to the numerical class label. Use :ocv:func:`CvMLData::get_class_labels_map` to obtain this map.
While reading the data, the method tries to define the type of variables (predictors and responses): ordered or categorical. If a value of the variable is not numerical (except for the label for a missing value), the type of the variable is set to ``CV_VAR_CATEGORICAL``. If all existing values of the variable are numerical, the type of the variable is set to ``CV_VAR_ORDERED``. So, the default definition of variables types works correctly for all cases except the case of a categorical variable with numerical class labels. In this case, the type ``CV_VAR_ORDERED`` is set. You should change the type to ``CV_VAR_CATEGORICAL`` using the method :ocv:func:`CvMLData::change_var_type`. For categorical variables, a common map is built to convert a string class label to the numerical class label. Use :ocv:func:`CvMLData::get_class_labels_map` to obtain this map.
Also, when reading the data, the method constructs the mask of missing values. For example, values are equal to `'?'`.
@@ -72,7 +72,7 @@ Returns a pointer to the matrix of predictors and response values
.. ocv:function:: const CvMat* CvMLData::get_values() const
The method returns a pointer to the matrix of predictor and response ``values`` or ``0`` if the data has not been loaded from the file yet.
The method returns a pointer to the matrix of predictor and response ``values`` or ``0`` if the data has not been loaded from the file yet.
The row count of this matrix equals the sample count. The column count equals predictors ``+ 1`` for the response (if exists) count. This means that each row of the matrix contains values of one sample predictor and response. The matrix type is ``CV_32FC1``.
@@ -82,7 +82,7 @@ Returns a pointer to the matrix of response values
.. ocv:function:: const CvMat* CvMLData::get_responses()
The method returns a pointer to the matrix of response values or throws an exception if the data has not been loaded from the file yet.
The method returns a pointer to the matrix of response values or throws an exception if the data has not been loaded from the file yet.
This is a single-column matrix of the type ``CV_32FC1``. Its row count is equal to the sample count, one column and .
@@ -92,7 +92,7 @@ Returns a pointer to the mask matrix of missing values
.. ocv:function:: const CvMat* CvMLData::get_missing() const
The method returns a pointer to the mask matrix of missing values or throws an exception if the data has not been loaded from the file yet.
The method returns a pointer to the mask matrix of missing values or throws an exception if the data has not been loaded from the file yet.
This matrix has the same size as the ``values`` matrix (see :ocv:func:`CvMLData::get_values`) and the type ``CV_8UC1``.
@@ -102,7 +102,7 @@ Specifies index of response column in the data matrix
.. ocv:function:: void CvMLData::set_response_idx( int idx )
The method sets the index of a response column in the ``values`` matrix (see :ocv:func:`CvMLData::get_values`) or throws an exception if the data has not been loaded from the file yet.
The method sets the index of a response column in the ``values`` matrix (see :ocv:func:`CvMLData::get_values`) or throws an exception if the data has not been loaded from the file yet.
The old response columns become predictors. If ``idx < 0``, there is no response.
@@ -115,15 +115,15 @@ Returns index of the response column in the loaded data matrix
The method returns the index of a response column in the ``values`` matrix (see :ocv:func:`CvMLData::get_values`) or throws an exception if the data has not been loaded from the file yet.
If ``idx < 0``, there is no response.
CvMLData::set_train_test_split
------------------------------
Divides the read data set into two disjoint training and test subsets.
Divides the read data set into two disjoint training and test subsets.
.. ocv:function:: void CvMLData::set_train_test_split( const CvTrainTestSplit * spl )
This method sets parameters for such a split using ``spl`` (see :ocv:class:`CvTrainTestSplit`) or throws an exception if the data has not been loaded from the file yet.
This method sets parameters for such a split using ``spl`` (see :ocv:class:`CvTrainTestSplit`) or throws an exception if the data has not been loaded from the file yet.
CvMLData::get_train_sample_idx
------------------------------
@@ -139,13 +139,13 @@ Returns the matrix of sample indices for a testing subset
.. ocv:function:: const CvMat* CvMLData::get_test_sample_idx() const
CvMLData::mix_train_and_test_idx
--------------------------------
Mixes the indices of training and test samples
.. ocv:function:: void CvMLData::mix_train_and_test_idx()
The method shuffles the indices of training and test samples preserving sizes of training and test subsets if the data split is set by :ocv:func:`CvMLData::get_values`. If the data has not been loaded from the file yet, an exception is thrown.
CvMLData::get_var_idx
@@ -153,8 +153,8 @@ CvMLData::get_var_idx
Returns the indices of the active variables in the data matrix
.. ocv:function:: const CvMat* CvMLData::get_var_idx()
The method returns the indices of variables (columns) used in the ``values`` matrix (see :ocv:func:`CvMLData::get_values`).
The method returns the indices of variables (columns) used in the ``values`` matrix (see :ocv:func:`CvMLData::get_values`).
It returns ``0`` if the used subset is not set. It throws an exception if the data has not been loaded from the file yet. Returned matrix is a single-row matrix of the type ``CV_32SC1``. Its column count is equal to the size of the used variable subset.
@@ -165,22 +165,22 @@ Enables or disables particular variable in the loaded data
.. ocv:function:: void CvMLData::change_var_idx( int vi, bool state )
By default, after reading the data set all variables in the ``values`` matrix (see :ocv:func:`CvMLData::get_values`) are used. But you may want to use only a subset of variables and include/exclude (depending on ``state`` value) a variable with the ``vi`` index from the used subset. If the data has not been loaded from the file yet, an exception is thrown.
CvMLData::get_var_types
-----------------------
Returns a matrix of the variable types.
Returns a matrix of the variable types.
.. ocv:function:: const CvMat* CvMLData::get_var_types()
The function returns a single-row matrix of the type ``CV_8UC1``, where each element is set to either ``CV_VAR_ORDERED`` or ``CV_VAR_CATEGORICAL``. The number of columns is equal to the number of variables. If data has not been loaded from file yet an exception is thrown.
CvMLData::set_var_types
-----------------------
Sets the variables types in the loaded data.
.. ocv:function:: void CvMLData::set_var_types( const char* str )
In the string, a variable type is followed by a list of variables indices. For example: ``"ord[0-17],cat[18]"``, ``"ord[0,2,4,10-12], cat[1,3,5-9,13,14]"``, ``"cat"`` (all variables are categorical), ``"ord"`` (all variables are ordered).
In the string, a variable type is followed by a list of variables indices. For example: ``"ord[0-17],cat[18]"``, ``"ord[0,2,4,10-12], cat[1,3,5-9,13,14]"``, ``"cat"`` (all variables are categorical), ``"ord"`` (all variables are ordered).
CvMLData::get_var_type
----------------------
@@ -189,15 +189,15 @@ Returns type of the specified variable
.. ocv:function:: int CvMLData::get_var_type( int var_idx ) const
The method returns the type of a variable by the index ``var_idx`` ( ``CV_VAR_ORDERED`` or ``CV_VAR_CATEGORICAL``).
CvMLData::change_var_type
-------------------------
Changes type of the specified variable
.. ocv:function:: void CvMLData::change_var_type( int var_idx, int type)
The method changes type of variable with index ``var_idx`` from existing type to ``type`` ( ``CV_VAR_ORDERED`` or ``CV_VAR_CATEGORICAL``).
CvMLData::set_delimiter
-----------------------
Sets the delimiter in the file used to separate input numbers
@@ -260,6 +260,6 @@ Structure setting the split of a data set read by :ocv:class:`CvMLData`.
There are two ways to construct a split:
* Set the training sample count (subset size) ``train_sample_count``. Other existing samples are located in a test subset.
* Set the training sample count (subset size) ``train_sample_count``. Other existing samples are located in a test subset.
* Set a training sample portion in ``[0,..1]``. The flag ``mix`` is used to mix training and test samples indices when the split is set. Otherwise, the data set is split in the storing order: the first part of samples of a given size is a training subset, the second part is a test subset.
+1 -1
View File
@@ -116,7 +116,7 @@ bool CvKNearest::train( const CvMat* _train_data, const CvMat* _responses,
if( !responses )
CV_ERROR( CV_StsNoMem, "Could not allocate memory for responses" );
if( _update_base && _dims != var_count )
CV_ERROR( CV_StsBadArg, "The newly added data have different dimensionality" );
+10 -10
View File
@@ -125,7 +125,7 @@ float icvCalcHaarPatternSum_2(
t -= read_sumTex( sumTex, sampler, (int2)(x + dx2.x, y + dy1.x), rows, cols, elemPerRow );
t += read_sumTex( sumTex, sampler, (int2)(x + dx2.x, y + dy2.x), rows, cols, elemPerRow );
d += t * src[4].x / ((dx2.x - dx1.x) * (dy2.x - dy1.x));
t = 0;
t += read_sumTex( sumTex, sampler, (int2)(x + dx1.y, y + dy1.y), rows, cols, elemPerRow );
t -= read_sumTex( sumTex, sampler, (int2)(x + dx1.y, y + dy2.y), rows, cols, elemPerRow );
@@ -161,14 +161,14 @@ float icvCalcHaarPatternSum_3(
t -= read_sumTex( sumTex, sampler, (int2)(x + dx2.x, y + dy1.x), rows, cols, elemPerRow );
t += read_sumTex( sumTex, sampler, (int2)(x + dx2.x, y + dy2.x), rows, cols, elemPerRow );
d += t * src[4].x / ((dx2.x - dx1.x) * (dy2.x - dy1.x));
t = 0;
t += read_sumTex( sumTex, sampler, (int2)(x + dx1.y, y + dy1.y), rows, cols, elemPerRow );
t -= read_sumTex( sumTex, sampler, (int2)(x + dx1.y, y + dy2.y), rows, cols, elemPerRow );
t -= read_sumTex( sumTex, sampler, (int2)(x + dx2.y, y + dy1.y), rows, cols, elemPerRow );
t += read_sumTex( sumTex, sampler, (int2)(x + dx2.y, y + dy2.y), rows, cols, elemPerRow );
d += t * src[4].y / ((dx2.y - dx1.y) * (dy2.y - dy1.y));
t = 0;
t += read_sumTex( sumTex, sampler, (int2)(x + dx1.z, y + dy1.z), rows, cols, elemPerRow );
t -= read_sumTex( sumTex, sampler, (int2)(x + dx1.z, y + dy2.z), rows, cols, elemPerRow );
@@ -204,21 +204,21 @@ float icvCalcHaarPatternSum_4(
t -= read_sumTex( sumTex, sampler, (int2)(x + dx2.x, y + dy1.x), rows, cols, elemPerRow );
t += read_sumTex( sumTex, sampler, (int2)(x + dx2.x, y + dy2.x), rows, cols, elemPerRow );
d += t * src[4].x / ((dx2.x - dx1.x) * (dy2.x - dy1.x));
t = 0;
t += read_sumTex( sumTex, sampler, (int2)(x + dx1.y, y + dy1.y), rows, cols, elemPerRow );
t -= read_sumTex( sumTex, sampler, (int2)(x + dx1.y, y + dy2.y), rows, cols, elemPerRow );
t -= read_sumTex( sumTex, sampler, (int2)(x + dx2.y, y + dy1.y), rows, cols, elemPerRow );
t += read_sumTex( sumTex, sampler, (int2)(x + dx2.y, y + dy2.y), rows, cols, elemPerRow );
d += t * src[4].y / ((dx2.y - dx1.y) * (dy2.y - dy1.y));
t = 0;
t += read_sumTex( sumTex, sampler, (int2)(x + dx1.z, y + dy1.z), rows, cols, elemPerRow );
t -= read_sumTex( sumTex, sampler, (int2)(x + dx1.z, y + dy2.z), rows, cols, elemPerRow );
t -= read_sumTex( sumTex, sampler, (int2)(x + dx2.z, y + dy1.z), rows, cols, elemPerRow );
t += read_sumTex( sumTex, sampler, (int2)(x + dx2.z, y + dy2.z), rows, cols, elemPerRow );
d += t * src[4].z / ((dx2.z - dx1.z) * (dy2.z - dy1.z));
t = 0;
t += read_sumTex( sumTex, sampler, (int2)(x + dx1.w, y + dy1.w), rows, cols, elemPerRow );
t -= read_sumTex( sumTex, sampler, (int2)(x + dx1.w, y + dy2.w), rows, cols, elemPerRow );
@@ -1231,7 +1231,7 @@ void compute_descriptors64(
barrier(CLK_LOCAL_MEM_FENCE);
reduce_sum25(sdx, sdy, sdxabs, sdyabs, tid);
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 25)
{
@@ -1361,14 +1361,14 @@ void reduce_sum128(volatile __local float* smem, int tid)
#if WAVE_SIZE < 64
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 32)
if (tid < 32)
{
#endif
smem[tid] += smem[tid + 32];
#if WAVE_SIZE < 32
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 16)
if (tid < 16)
{
#endif
smem[tid] += smem[tid + 16];
@@ -1415,7 +1415,7 @@ void reduce_sum64(volatile __local float* smem, int tid)
#if WAVE_SIZE < 32
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 16)
if (tid < 16)
{
#endif
smem[tid] += smem[tid + 16];
+1 -1
View File
@@ -53,7 +53,7 @@
#if defined(HAVE_OPENCV_GPU)
#include "opencv2/nonfree/gpu.hpp"
#if defined(HAVE_CUDA)
#include "opencv2/gpu/stream_accessor.hpp"
#include "opencv2/gpu/device/common.hpp"
+3 -3
View File
@@ -1964,10 +1964,10 @@ cvLoadHaarClassifierCascade( const char* directory, CvSize orig_window_size )
size += (n+1)*sizeof(char*);
const char** input_cascade = (const char**)cvAlloc( size );
if( !input_cascade )
CV_Error( CV_StsNoMem, "Could not allocate memory for input_cascade" );
char* ptr = (char*)(input_cascade + n + 1);
for( int i = 0; i < n; i++ )
@@ -1988,7 +1988,7 @@ cvLoadHaarClassifierCascade( const char* directory, CvSize orig_window_size )
}
input_cascade[n] = 0;
CvHaarClassifierCascade* cascade = icvLoadCascadeCART( input_cascade, n, orig_window_size );
if( input_cascade )
@@ -439,7 +439,7 @@ int CV_CascadeDetectorTest::detectMultiScale_C( const string& filename,
CvMat c_gray = grayImg;
CvSeq* rs = cvHaarDetectObjects(&c_gray, c_cascade, storage, 1.1, 3, flags[di] );
objects.clear();
for( int i = 0; i < rs->total; i++ )
{
+3 -3
View File
@@ -6,14 +6,14 @@ OpenCL Module Introduction
General Information
-------------------
The OpenCV OCL module contains a set of classes and functions that implement and accelerate select openCV functionality on OpenCL compatible devices. OpenCL is a Khronos standard, implemented by a variety of devices (CPUs, GPUs, FPGAs, ARM), abstracting the exact hardware details, while enabling vendors to provide native implementation for maximal acceleration on their hardware. The standard enjoys wide industry support, and the end user of the module will enjoy the data parallelism benefits that the specific platform/hardware may be capable of, in a platform/hardware independent manner.
The OpenCV OCL module contains a set of classes and functions that implement and accelerate select openCV functionality on OpenCL compatible devices. OpenCL is a Khronos standard, implemented by a variety of devices (CPUs, GPUs, FPGAs, ARM), abstracting the exact hardware details, while enabling vendors to provide native implementation for maximal acceleration on their hardware. The standard enjoys wide industry support, and the end user of the module will enjoy the data parallelism benefits that the specific platform/hardware may be capable of, in a platform/hardware independent manner.
While in the future we hope to validate (and enable) the OCL module in all OpenCL capable devices, we currently develop and test on GPU devices only. This includes both discrete GPUs (NVidia, AMD), as well as integrated chips(AMD APU and intel HD devices). Performance of any particular algorithm will depend on the particular platform characteristics and capabilities. However, currently (as of 2.4.4), accuracy and mathematical correctness has been verified to be identical to that of the pure CPU implementation on all tested GPU devices and platforms (both windows and linux).
While in the future we hope to validate (and enable) the OCL module in all OpenCL capable devices, we currently develop and test on GPU devices only. This includes both discrete GPUs (NVidia, AMD), as well as integrated chips(AMD APU and intel HD devices). Performance of any particular algorithm will depend on the particular platform characteristics and capabilities. However, currently (as of 2.4.4), accuracy and mathematical correctness has been verified to be identical to that of the pure CPU implementation on all tested GPU devices and platforms (both windows and linux).
The OpenCV OCL module includes utility functions, low-level vision primitives, and high-level algorithms. The utility functions and low-level primitives provide a powerful infrastructure for developing fast vision algorithms taking advangtage of OCL whereas the high-level functionality (samples)includes some state-of-the-art algorithms (including LK Optical flow, and Face detection) ready to be used by the application developers. The module is also accompanied by an extensive performance and accuracy test suite.
The OpenCV OCL module is designed for ease of use and does not require any knowledge of OpenCL. At a minimuml level, it can be viewed as a set of accelerators, that can take advantage of the high compute throughput that GPU/APU devices can provide. However, it can also be viewed as a starting point to really integratethe built-in functionality with your own custom OpenCL kernels, with or without modifying the source of OpenCV-OCL. Of course, knowledge of OpenCL will certainly help, however we hope that OpenCV-OCL module, and the kernels it contains in source code, can be very useful as a means of actually learning openCL. Such a knowledge would be necessary to further fine-tune any of the existing OpenCL kernels, or for extending the framework with new kernels. As of OpenCV 2.4.4, we introduce interoperability with OpenCL, enabling easy use of custom OpenCL kernels within the OpenCV framework.
The OpenCV OCL module is designed for ease of use and does not require any knowledge of OpenCL. At a minimuml level, it can be viewed as a set of accelerators, that can take advantage of the high compute throughput that GPU/APU devices can provide. However, it can also be viewed as a starting point to really integratethe built-in functionality with your own custom OpenCL kernels, with or without modifying the source of OpenCV-OCL. Of course, knowledge of OpenCL will certainly help, however we hope that OpenCV-OCL module, and the kernels it contains in source code, can be very useful as a means of actually learning openCL. Such a knowledge would be necessary to further fine-tune any of the existing OpenCL kernels, or for extending the framework with new kernels. As of OpenCV 2.4.4, we introduce interoperability with OpenCL, enabling easy use of custom OpenCL kernels within the OpenCV framework.
To use the OCL module, you need to make sure that you have the OpenCL SDK provided with your device vendor. To correctly run the OCL module, you need to have the OpenCL runtime provide by the device vendor, typically the device driver.
+3 -3
View File
@@ -500,13 +500,13 @@ Returns void
* **SORT_SELECTION** selection sort, currently cannot sort duplicate keys
* **SORT_MERGE** merge sort
* **SORT_RADIX** radix sort, only support signed int/float keys(``CV_32S``/``CV_32F``)
Returns the sorted result of all the elements in values based on equivalent keys.
The element unit in the values to be sorted is determined from the data type,
The element unit in the values to be sorted is determined from the data type,
i.e., a ``CV_32FC2`` input ``{a1a2, b1b2}`` will be considered as two elements, regardless its matrix dimension.
Both keys and values will be sorted inplace.
Both keys and values will be sorted inplace.
Keys needs to be a **single** channel `oclMat`.
+8 -8
View File
@@ -119,7 +119,7 @@ namespace cv
CV_EXPORTS void setDevice(Info &oclinfo, int devnum = 0);
//The two functions below enable other opencl program to use ocl module's cl_context and cl_command_queue
//returns cl_context *
//returns cl_context *
CV_EXPORTS void* getoclContext();
//returns cl_command_queue *
CV_EXPORTS void* getoclCommandQueue();
@@ -183,8 +183,8 @@ namespace cv
//! 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.
// 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
@@ -201,7 +201,7 @@ namespace cv
};
CV_EXPORTS void setBinaryDiskCache(int mode = CACHE_RELEASE, cv::String path = "./");
//! set where binary cache to be saved to
//! set where binary cache to be saved to
CV_EXPORTS void setBinpath(const char *path);
class CV_EXPORTS oclMatExpr;
@@ -514,10 +514,10 @@ 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
CV_EXPORTS Ptr<cv::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);
@@ -840,7 +840,7 @@ namespace cv
// supports CV_32FC1/CV_32FC2/CV_32FC4 data type
CV_EXPORTS void distanceToCenters(oclMat &dists, oclMat &labels, const oclMat &src, const oclMat &centers);
//!Does k-means procedure on GPU
//!Does k-means procedure on GPU
// supports CV_32FC1/CV_32FC2/CV_32FC4 data type
CV_EXPORTS double kmeans(const oclMat &src, int K, oclMat &bestLabels,
TermCriteria criteria, int attemps, int flags, oclMat &centers);
@@ -1683,7 +1683,7 @@ namespace cv
};
//! Returns the sorted result of all the elements in input based on equivalent keys.
//
// The element unit in the values to be sorted is determined from the data type,
// The element unit in the values to be sorted is determined from the data type,
// i.e., a CV_32FC2 input {a1a2, b1b2} will be considered as two elements, regardless its
// matrix dimension.
// both keys and values will be sorted inplace
@@ -130,7 +130,7 @@ namespace cv
{
openCLFree(tex_);
}
operator cl_mem()
operator cl_mem()
{
return tex_;
}
+1 -1
View File
@@ -99,7 +99,7 @@ int main(int argc, const char *argv[])
// set this to overwrite binary cache every time the test starts
ocl::setBinaryDiskCache(ocl::CACHE_UPDATE);
if (cmd.get<bool>("verify"))
{
TestSystem::instance().setNumIters(1);
+2 -2
View File
@@ -87,7 +87,7 @@ PERFTEST(StereoMatchBM)
d_bm(d_left, d_right, d_disp);
d_disp.download(disp);
GPU_FULL_OFF;
TestSystem::instance().setAccurate(-1, 0.);
}
@@ -98,4 +98,4 @@ PERFTEST(StereoMatchBM)
+2 -2
View File
@@ -284,7 +284,7 @@ PERFTEST(GaussianBlur)
Mat src, dst, ocl_dst;
int all_type[] = {CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4};
std::string type_name[] = {"CV_8UC1", "CV_8UC4", "CV_32FC1", "CV_32FC4"};
const int ksize = 7;
const int ksize = 7;
for (int size = Min_Size; size <= Max_Size; size *= Multiple)
{
@@ -374,4 +374,4 @@ PERFTEST(filter2D)
}
}
}
+2 -2
View File
@@ -55,14 +55,14 @@ PERFTEST(GoodFeaturesToTrack)
double qualityLevel = 0.01;
std::string images[] = { "rubberwhale1.png", "aloeL.jpg" };
std::vector<cv::Point2f> pts_gold, pts_ocl;
for(size_t imgIdx = 0; imgIdx < (sizeof(images)/sizeof(std::string)); ++imgIdx)
{
Mat frame = imread(abspath(images[imgIdx]), IMREAD_GRAYSCALE);
CV_Assert(!frame.empty());
for(float minDistance = 0; minDistance < 4; minDistance += 3.0)
{
SUBTEST << "image = " << images[imgIdx] << "; ";
+1 -1
View File
@@ -77,7 +77,7 @@ PERFTEST(HOG)
WARMUP_ON;
ocl_hog.detectMultiScale(d_src, d_found_locations);
WARMUP_OFF;
if(d_found_locations.size() == found_locations.size())
TestSystem::instance().setAccurate(1, 0);
else
+1 -1
View File
@@ -852,7 +852,7 @@ PERFTEST(meanShiftProc)
GPU_FULL_OFF;
vector<double> eps(2, 0.);
TestSystem::instance().ExpectMatsNear(dst, ocl_dst, eps);
TestSystem::instance().ExpectMatsNear(dst, ocl_dst, eps);
}
}
+1 -1
View File
@@ -72,7 +72,7 @@ PERFTEST(norm)
WARMUP_OFF;
d_src1.download(ocl_src1);
TestSystem::instance().ExpectedMatNear(src1, ocl_src1, .5);
TestSystem::instance().ExpectedMatNear(src1, ocl_src1, .5);
GPU_ON;
ocl::norm(d_src1, d_src2, NORM_INF);
+6 -6
View File
@@ -301,8 +301,8 @@ static const char* GetAnsiColorCode(GTestColor color) {
static void printMetricsUti(double cpu_time, double gpu_time, double gpu_full_time, double speedup, double fullspeedup, std::stringstream& stream, std::stringstream& cur_subtest_description)
{
//cout <<TAB<< setw(7) << stream.str();
cout <<TAB;
//cout <<TAB<< setw(7) << stream.str();
cout <<TAB;
stream.str("");
stream << cpu_time;
@@ -334,9 +334,9 @@ void TestSystem::printMetrics(int is_accurate, double cpu_time, double gpu_time,
stringstream stream;
std::stringstream &cur_subtest_description = getCurSubtestDescription();
#if GTEST_OS_WINDOWS&&!GTEST_OS_WINDOWS_MOBILE
WORD color;
const HANDLE stdout_handle = GetStdHandle(STD_OUTPUT_HANDLE);
// Gets the current text color.
@@ -401,10 +401,10 @@ void TestSystem::writeMetrics(double cpu_time, double gpu_time, double gpu_full_
exit(-1);
}
fprintf(record_, "%s,%s,%s,%.2f,%.3f,%.3f,%.3f,%.3f,%.3f,%.3f,%.3f,%.3f\n",
fprintf(record_, "%s,%s,%s,%.2f,%.3f,%.3f,%.3f,%.3f,%.3f,%.3f,%.3f,%.3f\n",
itname_changed_ ? itname_.c_str() : "",
cur_subtest_description_.str().c_str(),
_is_accurate_.c_str(),
_is_accurate_.c_str(),
accurate_diff_,
cpu_time, gpu_time, speedup, gpu_full_time, fullspeedup,
gpu_min, gpu_max, std_dev);
+1 -1
View File
@@ -133,7 +133,7 @@ PERFTEST(Split)
WARMUP_ON;
ocl::split(d_src, d_dst);
WARMUP_OFF;
WARMUP_OFF;
GPU_ON;
ocl::split(d_src, d_dst);
+3 -3
View File
@@ -118,7 +118,7 @@ inline int divUp(int total, int grain)
/////////////////////// add subtract multiply divide /////////////////////////
//////////////////////////////////////////////////////////////////////////////
template<typename T>
void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst,
void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst,
string kernelName, const char **kernelString, void *_scalar, int op_type = 0)
{
if(!src1.clCxt->supportsFeature(Context::CL_DOUBLE) && src1.type() == CV_64F)
@@ -187,12 +187,12 @@ void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst,
openCLExecuteKernel(clCxt, kernelString, kernelName, globalThreads, localThreads, args, -1, depth);
}
}
static void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst,
static void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst,
string kernelName, const char **kernelString, int op_type = 0)
{
arithmetic_run<char>(src1, src2, dst, kernelName, kernelString, (void *)NULL, op_type);
}
static void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst, const oclMat &mask,
static void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst, const oclMat &mask,
string kernelName, const char **kernelString, int op_type = 0)
{
if(!src1.clCxt->supportsFeature(Context::CL_DOUBLE) && src1.type() == CV_64F)
+8 -8
View File
@@ -218,7 +218,7 @@ public:
**Extend this if necessary later.
**Note that the kernel need to be further refined.
*/
static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
Size &ksize, const Point anchor, bool rectKernel)
{
//Normalize the result by default
@@ -275,8 +275,8 @@ static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
}
char compile_option[128];
sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D ERODE %s %s",
anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1],
sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D ERODE %s %s",
anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1],
s, rectKernel?"-D RECTKERNEL":"");
vector< pair<size_t, const void *> > args;
args.push_back(make_pair(sizeof(cl_mem), (void *)&src.data));
@@ -296,7 +296,7 @@ static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
//! data type supported: CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4
static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
Size &ksize, const Point anchor, bool rectKernel)
{
//Normalize the result by default
@@ -316,7 +316,7 @@ static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
Context *clCxt = src.clCxt;
string kernelName;
size_t localThreads[3] = {16, 16, 1};
size_t globalThreads[3] = {(src.cols + localThreads[0] - 1) / localThreads[0] *localThreads[0],
size_t globalThreads[3] = {(src.cols + localThreads[0] - 1) / localThreads[0] *localThreads[0],
(src.rows + localThreads[1] - 1) / localThreads[1] *localThreads[1], 1};
if (src.type() == CV_8UC1)
@@ -354,8 +354,8 @@ static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
}
char compile_option[128];
sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D DILATE %s %s",
anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1],
sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D DILATE %s %s",
anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1],
s, rectKernel?"-D RECTKERNEL":"");
vector< pair<size_t, const void *> > args;
args.push_back(make_pair(sizeof(cl_mem), (void *)&src.data));
@@ -622,7 +622,7 @@ static void GPUFilter2D(const oclMat &src, oclMat &dst, const oclMat &mat_kernel
int cn = src.oclchannels();
int src_step = (int)(src.step/src.elemSize());
int dst_step = (int)(dst.step/src.elemSize());
int localWidth = localThreads[0] + paddingPixels;
int localHeight = localThreads[1] + paddingPixels;
+11 -11
View File
@@ -147,7 +147,7 @@ struct Sorter<SELECTION>
static void sortCorners_caller(const EigType& eig_tex, oclMat& corners, const int count)
{
Context * cxt = Context::getContext();
size_t globalThreads[3] = {count, 1, 1};
size_t localThreads[3] = {GROUP_SIZE, 1, 1};
@@ -170,7 +170,7 @@ struct Sorter<SELECTION>
};
int findCorners_caller(
const TextureCL& eig,
const TextureCL& eig,
const float threshold,
const oclMat& mask,
oclMat& corners,
@@ -254,7 +254,7 @@ void cv::ocl::GoodFeaturesToTrackDetector_OCL::operator ()(const oclMat& image,
Sorter<SELECTION>::sortCorners_caller(*eig_tex, tmpCorners_, total);
}
}
if (minDistance < 1)
{
Rect roi_range(0, 0, maxCorners > 0 ? std::min(maxCorners, total) : total, 1);
@@ -338,14 +338,14 @@ void cv::ocl::GoodFeaturesToTrackDetector_OCL::downloadPoints(const oclMat &poin
CV_DbgAssert(points.type() == CV_32FC2);
points_v.resize(points.cols);
openCLSafeCall(clEnqueueReadBuffer(
*reinterpret_cast<cl_command_queue*>(getoclCommandQueue()),
reinterpret_cast<cl_mem>(points.data),
CL_TRUE,
0,
points.cols * sizeof(Point2f),
&points_v[0],
0,
NULL,
*reinterpret_cast<cl_command_queue*>(getoclCommandQueue()),
reinterpret_cast<cl_mem>(points.data),
CL_TRUE,
0,
points.cols * sizeof(Point2f),
&points_v[0],
0,
NULL,
NULL));
}
+2 -2
View File
@@ -1458,7 +1458,7 @@ void cv::ocl::OclCascadeClassifierBuf::CreateFactorRelatedBufs(
gimg1.release();
gsum.release();
gsqsum.release();
}
}
else if (!(m_flags & CV_HAAR_SCALE_IMAGE) && (flags & CV_HAAR_SCALE_IMAGE))
{
openCLSafeCall(clReleaseMemObject(((OclBuffers *)buffers)->newnodebuffer));
@@ -1476,7 +1476,7 @@ void cv::ocl::OclCascadeClassifierBuf::CreateFactorRelatedBufs(
{
return;
}
}
}
else
{
if (fabs(m_scaleFactor - scaleFactor) < 1e-6
+92 -92
View File
@@ -89,33 +89,33 @@ namespace cv
void compute_hists(int nbins, int block_stride_x, int blovck_stride_y,
int height, int width, const cv::ocl::oclMat &grad,
const cv::ocl::oclMat &qangle,
const cv::ocl::oclMat &qangle,
const cv::ocl::oclMat &gauss_w_lut, cv::ocl::oclMat &block_hists);
void normalize_hists(int nbins, int block_stride_x, int block_stride_y,
int height, int width, cv::ocl::oclMat &block_hists,
int height, int width, cv::ocl::oclMat &block_hists,
float threshold);
void classify_hists(int win_height, int win_width, int block_stride_y,
int block_stride_x, int win_stride_y, int win_stride_x,
int height, int width, const cv::ocl::oclMat &block_hists,
int block_stride_x, int win_stride_y, int win_stride_x,
int height, int width, const cv::ocl::oclMat &block_hists,
const cv::ocl::oclMat &coefs, float free_coef,
float threshold, cv::ocl::oclMat &labels);
void extract_descrs_by_rows(int win_height, int win_width, int block_stride_y,
int block_stride_x, int win_stride_y, int win_stride_x,
void extract_descrs_by_rows(int win_height, int win_width, int block_stride_y,
int block_stride_x, int win_stride_y, int win_stride_x,
int height, int width, const cv::ocl::oclMat &block_hists,
cv::ocl::oclMat &descriptors);
void extract_descrs_by_cols(int win_height, int win_width, int block_stride_y,
int block_stride_x, int win_stride_y, int win_stride_x,
void extract_descrs_by_cols(int win_height, int win_width, int block_stride_y,
int block_stride_x, int win_stride_y, int win_stride_x,
int height, int width, const cv::ocl::oclMat &block_hists,
cv::ocl::oclMat &descriptors);
void compute_gradients_8UC1(int height, int width, const cv::ocl::oclMat &img,
float angle_scale, cv::ocl::oclMat &grad,
float angle_scale, cv::ocl::oclMat &grad,
cv::ocl::oclMat &qangle, bool correct_gamma);
void compute_gradients_8UC4(int height, int width, const cv::ocl::oclMat &img,
float angle_scale, cv::ocl::oclMat &grad,
float angle_scale, cv::ocl::oclMat &grad,
cv::ocl::oclMat &qangle, bool correct_gamma);
}
}
@@ -129,8 +129,8 @@ static inline int divUp(int total, int grain)
return (total + grain - 1) / grain;
}
cv::ocl::HOGDescriptor::HOGDescriptor(Size win_size_, Size block_size_, Size block_stride_,
Size cell_size_, int nbins_, double win_sigma_,
cv::ocl::HOGDescriptor::HOGDescriptor(Size win_size_, Size block_size_, Size block_stride_,
Size cell_size_, int nbins_, double win_sigma_,
double threshold_L2hys_, bool gamma_correction_, int nlevels_)
: win_size(win_size_),
block_size(block_size_),
@@ -145,19 +145,19 @@ cv::ocl::HOGDescriptor::HOGDescriptor(Size win_size_, Size block_size_, Size blo
CV_Assert((win_size.width - block_size.width ) % block_stride.width == 0 &&
(win_size.height - block_size.height) % block_stride.height == 0);
CV_Assert(block_size.width % cell_size.width == 0 &&
CV_Assert(block_size.width % cell_size.width == 0 &&
block_size.height % cell_size.height == 0);
CV_Assert(block_stride == cell_size);
CV_Assert(cell_size == Size(8, 8));
Size cells_per_block(block_size.width / cell_size.width,
Size cells_per_block(block_size.width / cell_size.width,
block_size.height / cell_size.height);
CV_Assert(cells_per_block == Size(2, 2));
cv::Size blocks_per_win = numPartsWithin(win_size, block_size, block_stride);
hog::set_up_constants(nbins, block_stride.width, block_stride.height,
hog::set_up_constants(nbins, block_stride.width, block_stride.height,
blocks_per_win.width, blocks_per_win.height);
effect_size = Size(0, 0);
@@ -175,7 +175,7 @@ size_t cv::ocl::HOGDescriptor::getDescriptorSize() const
size_t cv::ocl::HOGDescriptor::getBlockHistogramSize() const
{
Size cells_per_block = Size(block_size.width / cell_size.width,
Size cells_per_block = Size(block_size.width / cell_size.width,
block_size.height / cell_size.height);
return (size_t)(nbins * cells_per_block.area());
}
@@ -189,7 +189,7 @@ bool cv::ocl::HOGDescriptor::checkDetectorSize() const
{
size_t detector_size = detector.rows * detector.cols;
size_t descriptor_size = getDescriptorSize();
return detector_size == 0 || detector_size == descriptor_size ||
return detector_size == 0 || detector_size == descriptor_size ||
detector_size == descriptor_size + 1;
}
@@ -230,7 +230,7 @@ void cv::ocl::HOGDescriptor::init_buffer(const oclMat &img, Size win_stride)
const size_t block_hist_size = getBlockHistogramSize();
const Size blocks_per_img = numPartsWithin(img.size(), block_size, block_stride);
block_hists.create(1,
block_hists.create(1,
static_cast<int>(block_hist_size * blocks_per_img.area()) + 256, CV_32F);
Size wins_per_img = numPartsWithin(img.size(), win_size, win_stride);
@@ -258,11 +258,11 @@ void cv::ocl::HOGDescriptor::computeGradient(const oclMat &img, oclMat &grad, oc
switch (img.type())
{
case CV_8UC1:
hog::compute_gradients_8UC1(effect_size.height, effect_size.width, img,
hog::compute_gradients_8UC1(effect_size.height, effect_size.width, img,
angleScale, grad, qangle, gamma_correction);
break;
case CV_8UC4:
hog::compute_gradients_8UC4(effect_size.height, effect_size.width, img,
hog::compute_gradients_8UC4(effect_size.height, effect_size.width, img,
angleScale, grad, qangle, gamma_correction);
break;
}
@@ -273,18 +273,18 @@ void cv::ocl::HOGDescriptor::computeBlockHistograms(const oclMat &img)
{
computeGradient(img, this->grad, this->qangle);
hog::compute_hists(nbins, block_stride.width, block_stride.height, effect_size.height,
hog::compute_hists(nbins, block_stride.width, block_stride.height, effect_size.height,
effect_size.width, grad, qangle, gauss_w_lut, block_hists);
hog::normalize_hists(nbins, block_stride.width, block_stride.height, effect_size.height,
hog::normalize_hists(nbins, block_stride.width, block_stride.height, effect_size.height,
effect_size.width, block_hists, (float)threshold_L2hys);
}
void cv::ocl::HOGDescriptor::getDescriptors(const oclMat &img, Size win_stride,
void cv::ocl::HOGDescriptor::getDescriptors(const oclMat &img, Size win_stride,
oclMat &descriptors, int descr_format)
{
CV_Assert(win_stride.width % block_stride.width == 0 &&
CV_Assert(win_stride.width % block_stride.width == 0 &&
win_stride.height % block_stride.height == 0);
init_buffer(img, win_stride);
@@ -295,19 +295,19 @@ void cv::ocl::HOGDescriptor::getDescriptors(const oclMat &img, Size win_stride,
Size blocks_per_win = numPartsWithin(win_size, block_size, block_stride);
Size wins_per_img = numPartsWithin(effect_size, win_size, win_stride);
descriptors.create(wins_per_img.area(),
descriptors.create(wins_per_img.area(),
static_cast<int>(blocks_per_win.area() * block_hist_size), CV_32F);
switch (descr_format)
{
case DESCR_FORMAT_ROW_BY_ROW:
hog::extract_descrs_by_rows(win_size.height, win_size.width,
block_stride.height, block_stride.width, win_stride.height, win_stride.width,
hog::extract_descrs_by_rows(win_size.height, win_size.width,
block_stride.height, block_stride.width, win_stride.height, win_stride.width,
effect_size.height, effect_size.width, block_hists, descriptors);
break;
case DESCR_FORMAT_COL_BY_COL:
hog::extract_descrs_by_cols(win_size.height, win_size.width,
block_stride.height, block_stride.width, win_stride.height, win_stride.width,
hog::extract_descrs_by_cols(win_size.height, win_size.width,
block_stride.height, block_stride.width, win_stride.height, win_stride.width,
effect_size.height, effect_size.width, block_hists, descriptors);
break;
default:
@@ -316,7 +316,7 @@ void cv::ocl::HOGDescriptor::getDescriptors(const oclMat &img, Size win_stride,
}
void cv::ocl::HOGDescriptor::detect(const oclMat &img, vector<Point> &hits,
void cv::ocl::HOGDescriptor::detect(const oclMat &img, vector<Point> &hits,
double hit_threshold, Size win_stride, Size padding)
{
CV_Assert(img.type() == CV_8UC1 || img.type() == CV_8UC4);
@@ -329,15 +329,15 @@ void cv::ocl::HOGDescriptor::detect(const oclMat &img, vector<Point> &hits,
if (win_stride == Size())
win_stride = block_stride;
else
CV_Assert(win_stride.width % block_stride.width == 0 &&
CV_Assert(win_stride.width % block_stride.width == 0 &&
win_stride.height % block_stride.height == 0);
init_buffer(img, win_stride);
computeBlockHistograms(img);
hog::classify_hists(win_size.height, win_size.width, block_stride.height,
block_stride.width, win_stride.height, win_stride.width,
effect_size.height, effect_size.width, block_hists, detector,
hog::classify_hists(win_size.height, win_size.width, block_stride.height,
block_stride.width, win_stride.height, win_stride.width,
effect_size.height, effect_size.width, block_hists, detector,
(float)free_coef, (float)hit_threshold, labels);
labels.download(labels_host);
@@ -354,8 +354,8 @@ void cv::ocl::HOGDescriptor::detect(const oclMat &img, vector<Point> &hits,
void cv::ocl::HOGDescriptor::detectMultiScale(const oclMat &img, vector<Rect> &found_locations,
double hit_threshold, Size win_stride, Size padding,
void cv::ocl::HOGDescriptor::detectMultiScale(const oclMat &img, vector<Rect> &found_locations,
double hit_threshold, Size win_stride, Size padding,
double scale0, int group_threshold)
{
CV_Assert(img.type() == CV_8UC1 || img.type() == CV_8UC4);
@@ -382,7 +382,7 @@ void cv::ocl::HOGDescriptor::detectMultiScale(const oclMat &img, vector<Rect> &f
if (win_stride == Size())
win_stride = block_stride;
else
CV_Assert(win_stride.width % block_stride.width == 0 &&
CV_Assert(win_stride.width % block_stride.width == 0 &&
win_stride.height % block_stride.height == 0);
init_buffer(img, win_stride);
image_scale.create(img.size(), img.type());
@@ -400,10 +400,10 @@ void cv::ocl::HOGDescriptor::detectMultiScale(const oclMat &img, vector<Rect> &f
resize(img, image_scale, effect_size);
detect(image_scale, locations, hit_threshold, win_stride, padding);
}
Size scaled_win_size(cvRound(win_size.width * scale),
Size scaled_win_size(cvRound(win_size.width * scale),
cvRound(win_size.height * scale));
for (size_t j = 0; j < locations.size(); j++)
all_candidates.push_back(Rect(Point2d((CvPoint)locations[j]) * scale,
all_candidates.push_back(Rect(Point2d((CvPoint)locations[j]) * scale,
scaled_win_size));
}
@@ -416,10 +416,10 @@ int cv::ocl::HOGDescriptor::numPartsWithin(int size, int part_size, int stride)
return (size - part_size + stride) / stride;
}
cv::Size cv::ocl::HOGDescriptor::numPartsWithin(cv::Size size, cv::Size part_size,
cv::Size cv::ocl::HOGDescriptor::numPartsWithin(cv::Size size, cv::Size part_size,
cv::Size stride)
{
return Size(numPartsWithin(size.width, part_size.width, stride.width),
return Size(numPartsWithin(size.width, part_size.width, stride.width),
numPartsWithin(size.height, part_size.height, stride.height));
}
@@ -1601,8 +1601,8 @@ static int power_2up(unsigned int n)
return -1; // Input is too big
}
void cv::ocl::device::hog::set_up_constants(int nbins,
int block_stride_x, int block_stride_y,
void cv::ocl::device::hog::set_up_constants(int nbins,
int block_stride_x, int block_stride_y,
int nblocks_win_x, int nblocks_win_y)
{
cnbins = nbins;
@@ -1622,21 +1622,21 @@ void cv::ocl::device::hog::set_up_constants(int nbins,
cdescr_size = descr_size;
}
void cv::ocl::device::hog::compute_hists(int nbins,
void cv::ocl::device::hog::compute_hists(int nbins,
int block_stride_x, int block_stride_y,
int height, int width,
const cv::ocl::oclMat &grad,
const cv::ocl::oclMat &qangle,
const cv::ocl::oclMat &gauss_w_lut,
int height, int width,
const cv::ocl::oclMat &grad,
const cv::ocl::oclMat &qangle,
const cv::ocl::oclMat &gauss_w_lut,
cv::ocl::oclMat &block_hists)
{
Context *clCxt = Context::getContext();
vector< pair<size_t, const void *> > args;
string kernelName = "compute_hists_lut_kernel";
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x)
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x)
/ block_stride_x;
int img_block_height = (height - CELLS_PER_BLOCK_Y * CELL_HEIGHT + block_stride_y)
int img_block_height = (height - CELLS_PER_BLOCK_Y * CELL_HEIGHT + block_stride_y)
/ block_stride_y;
int blocks_total = img_block_width * img_block_height;
@@ -1645,7 +1645,7 @@ void cv::ocl::device::hog::compute_hists(int nbins,
int blocks_in_group = 4;
size_t localThreads[3] = { blocks_in_group * 24, 2, 1 };
size_t globalThreads[3] = {
size_t globalThreads[3] = {
divUp(img_block_width * img_block_height, blocks_in_group) * localThreads[0], 2, 1 };
int hists_size = (nbins * CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y * 12) * sizeof(float);
@@ -1669,7 +1669,7 @@ void cv::ocl::device::hog::compute_hists(int nbins,
if(hog_device_cpu)
{
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
localThreads, args, -1, -1, "-D CPU");
}else
{
@@ -1677,15 +1677,15 @@ void cv::ocl::device::hog::compute_hists(int nbins,
int wave_size = queryDeviceInfo<WAVEFRONT_SIZE, int>(kernel);
char opt[32] = {0};
sprintf(opt, "-D WAVE_SIZE=%d", wave_size);
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
localThreads, args, -1, -1, opt);
}
}
void cv::ocl::device::hog::normalize_hists(int nbins,
void cv::ocl::device::hog::normalize_hists(int nbins,
int block_stride_x, int block_stride_y,
int height, int width,
cv::ocl::oclMat &block_hists,
int height, int width,
cv::ocl::oclMat &block_hists,
float threshold)
{
Context *clCxt = Context::getContext();
@@ -1693,14 +1693,14 @@ void cv::ocl::device::hog::normalize_hists(int nbins,
string kernelName;
int block_hist_size = nbins * CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y;
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x)
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x)
/ block_stride_x;
int img_block_height = (height - CELLS_PER_BLOCK_Y * CELL_HEIGHT + block_stride_y)
int img_block_height = (height - CELLS_PER_BLOCK_Y * CELL_HEIGHT + block_stride_y)
/ block_stride_y;
int nthreads;
size_t globalThreads[3] = { 1, 1, 1 };
size_t localThreads[3] = { 1, 1, 1 };
if ( nbins == 9 )
{
/* optimized for the case of 9 bins */
@@ -1720,7 +1720,7 @@ void cv::ocl::device::hog::normalize_hists(int nbins,
localThreads[0] = nthreads;
if ((nthreads < 32) || (nthreads > 512) )
cv::ocl::error("normalize_hists: histogram's size is too small or too big",
cv::ocl::error("normalize_hists: histogram's size is too small or too big",
__FILE__, __LINE__, "normalize_hists");
args.push_back( make_pair( sizeof(cl_int), (void *)&nthreads));
@@ -1733,7 +1733,7 @@ void cv::ocl::device::hog::normalize_hists(int nbins,
args.push_back( make_pair( nthreads * sizeof(float), (void *)NULL));
if(hog_device_cpu)
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
localThreads, args, -1, -1, "-D CPU");
else
{
@@ -1741,18 +1741,18 @@ void cv::ocl::device::hog::normalize_hists(int nbins,
int wave_size = queryDeviceInfo<WAVEFRONT_SIZE, int>(kernel);
char opt[32] = {0};
sprintf(opt, "-D WAVE_SIZE=%d", wave_size);
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
localThreads, args, -1, -1, opt);
}
}
void cv::ocl::device::hog::classify_hists(int win_height, int win_width,
int block_stride_y, int block_stride_x,
int win_stride_y, int win_stride_x,
int height, int width,
const cv::ocl::oclMat &block_hists,
const cv::ocl::oclMat &coefs,
float free_coef, float threshold,
void cv::ocl::device::hog::classify_hists(int win_height, int win_width,
int block_stride_y, int block_stride_x,
int win_stride_y, int win_stride_x,
int height, int width,
const cv::ocl::oclMat &block_hists,
const cv::ocl::oclMat &coefs,
float free_coef, float threshold,
cv::ocl::oclMat &labels)
{
Context *clCxt = Context::getContext();
@@ -1785,7 +1785,7 @@ void cv::ocl::device::hog::classify_hists(int win_height, int win_width,
int win_block_stride_y = win_stride_y / block_stride_y;
int img_win_width = (width - win_width + win_stride_x) / win_stride_x;
int img_win_height = (height - win_height + win_stride_y) / win_stride_y;
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) /
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) /
block_stride_x;
size_t globalThreads[3] = { img_win_width * nthreads, img_win_height, 1 };
@@ -1802,7 +1802,7 @@ void cv::ocl::device::hog::classify_hists(int win_height, int win_width,
args.push_back( make_pair( sizeof(cl_mem), (void *)&labels.data));
if(hog_device_cpu)
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
localThreads, args, -1, -1, "-D CPU");
else
{
@@ -1810,16 +1810,16 @@ void cv::ocl::device::hog::classify_hists(int win_height, int win_width,
int wave_size = queryDeviceInfo<WAVEFRONT_SIZE, int>(kernel);
char opt[32] = {0};
sprintf(opt, "-D WAVE_SIZE=%d", wave_size);
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
localThreads, args, -1, -1, opt);
}
}
void cv::ocl::device::hog::extract_descrs_by_rows(int win_height, int win_width,
void cv::ocl::device::hog::extract_descrs_by_rows(int win_height, int win_width,
int block_stride_y, int block_stride_x,
int win_stride_y, int win_stride_x,
int win_stride_y, int win_stride_x,
int height, int width,
const cv::ocl::oclMat &block_hists,
const cv::ocl::oclMat &block_hists,
cv::ocl::oclMat &descriptors)
{
Context *clCxt = Context::getContext();
@@ -1830,7 +1830,7 @@ void cv::ocl::device::hog::extract_descrs_by_rows(int win_height, int win_width,
int win_block_stride_y = win_stride_y / block_stride_y;
int img_win_width = (width - win_width + win_stride_x) / win_stride_x;
int img_win_height = (height - win_height + win_stride_y) / win_stride_y;
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) /
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) /
block_stride_x;
int descriptors_quadstep = descriptors.step >> 2;
@@ -1847,15 +1847,15 @@ void cv::ocl::device::hog::extract_descrs_by_rows(int win_height, int win_width,
args.push_back( make_pair( sizeof(cl_mem), (void *)&block_hists.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&descriptors.data));
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
localThreads, args, -1, -1);
}
void cv::ocl::device::hog::extract_descrs_by_cols(int win_height, int win_width,
void cv::ocl::device::hog::extract_descrs_by_cols(int win_height, int win_width,
int block_stride_y, int block_stride_x,
int win_stride_y, int win_stride_x,
int win_stride_y, int win_stride_x,
int height, int width,
const cv::ocl::oclMat &block_hists,
const cv::ocl::oclMat &block_hists,
cv::ocl::oclMat &descriptors)
{
Context *clCxt = Context::getContext();
@@ -1866,7 +1866,7 @@ void cv::ocl::device::hog::extract_descrs_by_cols(int win_height, int win_width,
int win_block_stride_y = win_stride_y / block_stride_y;
int img_win_width = (width - win_width + win_stride_x) / win_stride_x;
int img_win_height = (height - win_height + win_stride_y) / win_stride_y;
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) /
int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) /
block_stride_x;
int descriptors_quadstep = descriptors.step >> 2;
@@ -1884,15 +1884,15 @@ void cv::ocl::device::hog::extract_descrs_by_cols(int win_height, int win_width,
args.push_back( make_pair( sizeof(cl_mem), (void *)&block_hists.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&descriptors.data));
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
localThreads, args, -1, -1);
}
void cv::ocl::device::hog::compute_gradients_8UC1(int height, int width,
void cv::ocl::device::hog::compute_gradients_8UC1(int height, int width,
const cv::ocl::oclMat &img,
float angle_scale,
cv::ocl::oclMat &grad,
cv::ocl::oclMat &qangle,
float angle_scale,
cv::ocl::oclMat &grad,
cv::ocl::oclMat &qangle,
bool correct_gamma)
{
Context *clCxt = Context::getContext();
@@ -1918,15 +1918,15 @@ void cv::ocl::device::hog::compute_gradients_8UC1(int height, int width,
args.push_back( make_pair( sizeof(cl_char), (void *)&correctGamma));
args.push_back( make_pair( sizeof(cl_int), (void *)&cnbins));
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
localThreads, args, -1, -1);
}
void cv::ocl::device::hog::compute_gradients_8UC4(int height, int width,
void cv::ocl::device::hog::compute_gradients_8UC4(int height, int width,
const cv::ocl::oclMat &img,
float angle_scale,
cv::ocl::oclMat &grad,
cv::ocl::oclMat &qangle,
float angle_scale,
cv::ocl::oclMat &grad,
cv::ocl::oclMat &qangle,
bool correct_gamma)
{
Context *clCxt = Context::getContext();
@@ -1953,6 +1953,6 @@ void cv::ocl::device::hog::compute_gradients_8UC4(int height, int width,
args.push_back( make_pair( sizeof(cl_char), (void *)&correctGamma));
args.push_back( make_pair( sizeof(cl_int), (void *)&cnbins));
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads,
localThreads, args, -1, -1);
}
+3 -3
View File
@@ -291,7 +291,7 @@ namespace cv
args.push_back( make_pair(sizeof(cl_int), (void *)&map1.cols));
args.push_back( make_pair(sizeof(cl_int), (void *)&map1.rows));
args.push_back( make_pair(sizeof(cl_int), (void *)&cols));
if(src.clCxt->supportsFeature(Context::CL_DOUBLE))
{
args.push_back( make_pair(sizeof(cl_double4), (void *)&borderValue));
@@ -1115,7 +1115,7 @@ namespace cv
args.push_back( make_pair( sizeof(cl_int) , (void *)&sum.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&sum_offset));
size_t gt2[3] = {t_sum.cols * 32, 1, 1}, lt2[3] = {256, 1, 1};
openCLExecuteKernel(src.clCxt, &imgproc_integral_sum, "integral_sum_rows", gt2, lt2, args, -1, depth);
openCLExecuteKernel(src.clCxt, &imgproc_integral_sum, "integral_sum_rows", gt2, lt2, args, -1, depth);
}
/////////////////////// corner //////////////////////////////
@@ -1230,7 +1230,7 @@ namespace cv
oclMat dx, dy;
cornerMinEigenVal_dxdy(src, dst, dx, dy, blockSize, ksize, borderType);
}
void cornerMinEigenVal_dxdy(const oclMat &src, oclMat &dst, oclMat &dx, oclMat &dy, int blockSize, int ksize, int borderType)
{
if(!src.clCxt->supportsFeature(Context::CL_DOUBLE) && src.depth() == CV_64F)
+5 -5
View File
@@ -168,7 +168,7 @@ namespace cv
};
// global variables to hold binary cache properties
static int enable_disk_cache =
static int enable_disk_cache =
#ifdef _DEBUG
false;
#else
@@ -514,8 +514,8 @@ namespace cv
return;
}
update_disk_cache |= (mode & CACHE_UPDATE) == CACHE_UPDATE;
enable_disk_cache |=
#ifdef _DEBUG
enable_disk_cache |=
#ifdef _DEBUG
(mode & CACHE_DEBUG) == CACHE_DEBUG;
#else
(mode & CACHE_RELEASE) == CACHE_RELEASE;
@@ -950,8 +950,8 @@ namespace cv
bool initialized()
{
return *((volatile int*)&Context::val) != 0 &&
Context::clCxt->impl->clCmdQueue != NULL&&
return *((volatile int*)&Context::val) != 0 &&
Context::clCxt->impl->clCmdQueue != NULL&&
Context::clCxt->impl->oclcontext != NULL;
}
+3 -3
View File
@@ -103,7 +103,7 @@ namespace cv
{
// FIXME!
// always use naive until convolve is imported
return true;
return true;
}
//////////////////////////////////////////////////////////////////////
@@ -120,7 +120,7 @@ namespace cv
else
{
buf.image_sqsums.resize(1);
// TODO, add double support for ocl::integral
// use CPU integral temporarily
Mat sums, sqsums;
@@ -360,7 +360,7 @@ namespace cv
}
else
{
split(image, buf.images);
templ_sum = sum(templ) / templ.size().area();
buf.image_sums.resize(buf.images.size());
+1 -1
View File
@@ -627,7 +627,7 @@ static void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, stri
if(Context::getContext()->supportsFeature(Context::CL_VER_1_2) &&
dst.offset == 0 && dst.cols == dst.wholecols)
{
clEnqueueFillBuffer((cl_command_queue)dst.clCxt->oclCommandQueue(),
clEnqueueFillBuffer((cl_command_queue)dst.clCxt->oclCommandQueue(),
(cl_mem)dst.data, args[0].second, args[0].first, 0, dst.step * dst.rows, 0, NULL, NULL);
}
else
+2 -2
View File
@@ -197,7 +197,7 @@ namespace cv
desc.buffer = NULL;
desc.num_mip_levels = 0;
desc.num_samples = 0;
texture = clCreateImage((cl_context)mat.clCxt->oclContext(), CL_MEM_READ_WRITE, &format, &desc, NULL, &err);
texture = clCreateImage((cl_context)mat.clCxt->oclContext(), CL_MEM_READ_WRITE, &format, &desc, NULL, &err);
}
else
#endif
@@ -223,7 +223,7 @@ namespace cv
const size_t regin[3] = {mat.cols * mat.elemSize(), mat.rows, 1};
clEnqueueCopyBufferRect((cl_command_queue)mat.clCxt->oclCommandQueue(), (cl_mem)mat.data, devData, origin, origin,
regin, mat.step, 0, mat.cols * mat.elemSize(), 0, 0, NULL, NULL);
clFlush((cl_command_queue)mat.clCxt->oclCommandQueue());
clFlush((cl_command_queue)mat.clCxt->oclCommandQueue());
}
else
{
+1 -1
View File
@@ -143,7 +143,7 @@ static void icvContourMoments( CvSeq* contour, CvMoments* mom )
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step ));
openCLExecuteKernel(dst_a.clCxt, &moments, "icvContourMoments", globalThreads, localThreads, args, -1, -1);
cv::Mat dst(dst_a);
a00 = a10 = a01 = a20 = a11 = a02 = a30 = a21 = a12 = a03 = 0.0;
if (!cv::ocl::Context::getContext()->supportsFeature(Context::CL_DOUBLE))
+10 -10
View File
@@ -66,7 +66,7 @@ __kernel void arithm_absdiff_D0 (__global uchar *src1, int src1_step, int src1_o
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -117,7 +117,7 @@ __kernel void arithm_absdiff_D2 (__global ushort *src1, int src1_step, int src1_
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -154,7 +154,7 @@ __kernel void arithm_absdiff_D3 (__global short *src1, int src1_step, int src1_o
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -261,7 +261,7 @@ __kernel void arithm_s_absdiff_C1_D0 (__global uchar *src1, int src1_step, int
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -303,7 +303,7 @@ __kernel void arithm_s_absdiff_C1_D2 (__global ushort *src1, int src1_step, in
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -337,7 +337,7 @@ __kernel void arithm_s_absdiff_C1_D3 (__global short *src1, int src1_step, int
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -443,7 +443,7 @@ __kernel void arithm_s_absdiff_C2_D0 (__global uchar *src1, int src1_step, int
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -588,7 +588,7 @@ __kernel void arithm_s_absdiff_C3_D0 (__global uchar *src1, int src1_step, int
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -645,7 +645,7 @@ __kernel void arithm_s_absdiff_C3_D2 (__global ushort *src1, int src1_step, in
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -698,7 +698,7 @@ __kernel void arithm_s_absdiff_C3_D3 (__global short *src1, int src1_step, int
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
+5 -5
View File
@@ -128,7 +128,7 @@ __kernel void addWeighted_D2 (__global ushort *src1, int src1_step,int src1_offs
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -191,7 +191,7 @@ __kernel void addWeighted_D3 (__global short *src1, int src1_step,int src1_offse
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -255,7 +255,7 @@ __kernel void addWeighted_D4 (__global int *src1, int src1_step,int src1_offset,
x = x << 2;
#define bitOfInt (sizeof(int)== 4 ? 2: 3)
#ifdef dst_align
#undef dst_align
#endif
@@ -319,7 +319,7 @@ __kernel void addWeighted_D5 (__global float *src1,int src1_step,int src1_offset
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -384,7 +384,7 @@ __kernel void addWeighted_D6 (__global double *src1, int src1_step,int src1_offs
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
+4 -4
View File
@@ -67,7 +67,7 @@ __kernel void arithm_s_add_C1_D0 (__global uchar *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -110,7 +110,7 @@ __kernel void arithm_s_add_C1_D2 (__global ushort *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -145,7 +145,7 @@ __kernel void arithm_s_add_C1_D3 (__global short *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -250,7 +250,7 @@ __kernel void arithm_s_add_C2_D0 (__global uchar *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -69,7 +69,7 @@ __kernel void arithm_s_add_with_mask_C1_D0 (__global uchar *src1, int src1_ste
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -122,7 +122,7 @@ __kernel void arithm_s_add_with_mask_C1_D2 (__global ushort *src1, int src1_st
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -160,7 +160,7 @@ __kernel void arithm_s_add_with_mask_C1_D3 (__global short *src1, int src1_ste
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -284,7 +284,7 @@ __kernel void arithm_s_add_with_mask_C2_D0 (__global uchar *src1, int src1_ste
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -72,7 +72,7 @@ __kernel void arithm_bitwise_binary_D0 (__global uchar *src1, int src1_step, int
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -125,7 +125,7 @@ __kernel void arithm_bitwise_binary_D1 (__global char *src1, int src1_step, int
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -179,7 +179,7 @@ __kernel void arithm_bitwise_binary_D2 (__global ushort *src1, int src1_step, in
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -234,7 +234,7 @@ __kernel void arithm_bitwise_binary_D3 (__global short *src1, int src1_step, int
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -73,7 +73,7 @@ __kernel void arithm_bitwise_binary_with_mask_C1_D0 (
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -118,7 +118,7 @@ __kernel void arithm_bitwise_binary_with_mask_C1_D1 (
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -163,7 +163,7 @@ __kernel void arithm_bitwise_binary_with_mask_C1_D2 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -206,7 +206,7 @@ __kernel void arithm_bitwise_binary_with_mask_C1_D3 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -349,7 +349,7 @@ __kernel void arithm_bitwise_binary_with_mask_C2_D0 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -391,7 +391,7 @@ __kernel void arithm_bitwise_binary_with_mask_C2_D1 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -70,7 +70,7 @@ __kernel void arithm_s_bitwise_binary_C1_D0 (
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -108,7 +108,7 @@ __kernel void arithm_s_bitwise_binary_C1_D1 (
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -146,7 +146,7 @@ __kernel void arithm_s_bitwise_binary_C1_D2 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -181,7 +181,7 @@ __kernel void arithm_s_bitwise_binary_C1_D3 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -293,7 +293,7 @@ __kernel void arithm_s_bitwise_binary_C2_D0 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -331,7 +331,7 @@ __kernel void arithm_s_bitwise_binary_C2_D1 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -71,7 +71,7 @@ __kernel void arithm_s_bitwise_binary_with_mask_C1_D0 (
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -113,7 +113,7 @@ __kernel void arithm_s_bitwise_binary_with_mask_C1_D1 (
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -154,7 +154,7 @@ __kernel void arithm_s_bitwise_binary_with_mask_C1_D2 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -192,7 +192,7 @@ __kernel void arithm_s_bitwise_binary_with_mask_C1_D3 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -318,7 +318,7 @@ __kernel void arithm_s_bitwise_binary_with_mask_C2_D0 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
@@ -358,7 +358,7 @@ __kernel void arithm_s_bitwise_binary_with_mask_C2_D1 (
if (x < cols && y < rows)
{
x = x << 1;
#ifdef dst_align
#undef dst_align
#endif
+4 -4
View File
@@ -63,7 +63,7 @@ __kernel void arithm_bitwise_not_D0 (__global uchar *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -106,7 +106,7 @@ __kernel void arithm_bitwise_not_D1 (__global char *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -143,7 +143,7 @@ __kernel void arithm_bitwise_not_D2 (__global ushort *src1, int src1_step, int s
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -181,7 +181,7 @@ __kernel void arithm_bitwise_not_D3 (__global short *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
+12 -12
View File
@@ -65,7 +65,7 @@ __kernel void arithm_compare_eq_D0 (__global uchar *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -120,7 +120,7 @@ __kernel void arithm_compare_ne_D2 (__global ushort *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -174,7 +174,7 @@ __kernel void arithm_compare_eq_D3 (__global short *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -380,7 +380,7 @@ __kernel void arithm_compare_gt_D0 (__global uchar *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -434,7 +434,7 @@ __kernel void arithm_compare_gt_D2 (__global ushort *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -490,7 +490,7 @@ __kernel void arithm_compare_gt_D3 (__global short *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -700,7 +700,7 @@ __kernel void arithm_compare_ge_D0 (__global uchar *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -757,7 +757,7 @@ __kernel void arithm_compare_ge_D2 (__global ushort *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -815,7 +815,7 @@ __kernel void arithm_compare_ge_D3 (__global short *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -869,7 +869,7 @@ __kernel void arithm_compare_ge_D4 (__global int *src1, int src1_step, int src1_
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -921,7 +921,7 @@ __kernel void arithm_compare_ge_D5 (__global float *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -975,7 +975,7 @@ __kernel void arithm_compare_ge_D6 (__global double *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
+9 -9
View File
@@ -61,7 +61,7 @@ __kernel void arithm_compare_ne_D0 (__global uchar *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -116,7 +116,7 @@ __kernel void arithm_compare_ne_D2 (__global ushort *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -171,7 +171,7 @@ __kernel void arithm_compare_ne_D3 (__global short *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -380,7 +380,7 @@ __kernel void arithm_compare_lt_D0 (__global uchar *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -435,7 +435,7 @@ __kernel void arithm_compare_lt_D2 (__global ushort *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -491,7 +491,7 @@ __kernel void arithm_compare_lt_D3 (__global short *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -704,7 +704,7 @@ __kernel void arithm_compare_le_D0 (__global uchar *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -760,7 +760,7 @@ __kernel void arithm_compare_le_D2 (__global ushort *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -816,7 +816,7 @@ __kernel void arithm_compare_le_D3 (__global short *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
+6 -6
View File
@@ -93,7 +93,7 @@ __kernel void arithm_div_D0 (__global uchar *src1, int src1_step, int src1_offse
if (coor.x < cols && coor.y < rows)
{
coor.x = coor.x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -137,7 +137,7 @@ __kernel void arithm_div_D2 (__global ushort *src1, int src1_step, int src1_offs
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -180,7 +180,7 @@ __kernel void arithm_div_D3 (__global short *src1, int src1_step, int src1_offse
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -298,7 +298,7 @@ __kernel void arithm_s_div_D0 (__global uchar *src, int src_step, int src_offset
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -337,7 +337,7 @@ __kernel void arithm_s_div_D2 (__global ushort *src, int src_step, int src_offse
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -375,7 +375,7 @@ __kernel void arithm_s_div_D3 (__global short *src, int src_step, int src_offset
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
+4 -4
View File
@@ -64,7 +64,7 @@ __kernel void arithm_flip_rows_D0 (__global uchar *src, int src_step, int src_of
if (x < cols && y < thread_rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -122,7 +122,7 @@ __kernel void arithm_flip_rows_D1 (__global char *src, int src_step, int src_off
if (x < cols && y < thread_rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -167,7 +167,7 @@ __kernel void arithm_flip_rows_D2 (__global ushort *src, int src_step, int src_o
if (x < cols && y < thread_rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -212,7 +212,7 @@ __kernel void arithm_flip_rows_D3 (__global short *src, int src_step, int src_of
if (x < cols && y < thread_rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
+3 -3
View File
@@ -90,7 +90,7 @@ __kernel void arithm_mul_D0 (__global uchar *src1, int src1_step, int src1_offse
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -138,7 +138,7 @@ __kernel void arithm_mul_D2 (__global ushort *src1, int src1_step, int src1_offs
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
@@ -177,7 +177,7 @@ __kernel void arithm_mul_D3 (__global short *src1, int src1_step, int src1_offse
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
+1 -1
View File
@@ -425,7 +425,7 @@ __kernel void BruteForceMatch_RadiusMatch(
barrier(CLK_LOCAL_MEM_FENCE);
}
if (queryIdx < query_rows && trainIdx < train_rows &&
if (queryIdx < query_rows && trainIdx < train_rows &&
convert_float(result) < maxDistance/* && mask(queryIdx, trainIdx)*/)
{
unsigned int ind = atom_inc(nMatches + queryIdx);
+1 -1
View File
@@ -465,7 +465,7 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
start_addr = mad24(y,dst_step_in_pixel,x);
dst[start_addr] = sum;
}
}
@@ -231,7 +231,7 @@ __kernel void boxFilter_C1_D0(__global const uchar * restrict src, __global ucha
{
tmp_sum += (data[i]);
}
int index = dst_startY * dst_step + dst_startX + (col-anX)*4;
temp[0][col] = tmp_sum + (data[0]);
+2 -2
View File
@@ -46,7 +46,7 @@
typedef int sumtype;
typedef float sqsumtype;
#ifndef STUMP_BASED
#ifndef STUMP_BASED
#define STUMP_BASED 1
#endif
@@ -323,7 +323,7 @@ __kernel void __attribute__((reqd_work_group_size(8,8,1)))gpuRunHaarClassifierCa
int root_offset = 0;
for(int lcl_loop=0; lcl_loop<lcl_loops && tempnodecounter<stageinfo.x;)
{
__global GpuHidHaarTreeNode* currentnodeptr =
__global GpuHidHaarTreeNode* currentnodeptr =
nodeptr + (nodecounter + tempnodecounter) * stump_factor + root_offset;
int4 info1 = *(__global int4*)(&(currentnodeptr->p[0][0]));
@@ -207,7 +207,7 @@ __kernel void gpuRunHaarClassifierCascade_scaled2(
- sum[clamp(mad24(info3.y, step, info3.z), 0, max_idx)] -
sum[clamp(mad24(info3.w, step, info3.x), 0, max_idx)]
+ sum[clamp(mad24(info3.w, step, info3.z), 0, max_idx)]) * w.z;
bool passThres = classsum >= nodethreshold;
#if STUMP_BASED
+8 -8
View File
@@ -71,15 +71,15 @@ void reduce(volatile __local int* smem, int val, int tid)
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)
@@ -125,15 +125,15 @@ void reduce(__local volatile int* smem, int val, int tid)
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)
+16 -16
View File
@@ -49,12 +49,12 @@
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
inline float ELEM_INT2(image2d_t _eig, int _x, int _y)
inline float ELEM_INT2(image2d_t _eig, int _x, int _y)
{
return read_imagef(_eig, sampler, (int2)(_x, _y)).x;
}
inline float ELEM_FLT2(image2d_t _eig, float2 pt)
inline float ELEM_FLT2(image2d_t _eig, float2 pt)
{
return read_imagef(_eig, sampler, pt).x;
}
@@ -132,7 +132,7 @@ __kernel
const int pairDistance = 1 << (stage - passOfStage);
const int blockWidth = 2 * pairDistance;
const int leftId = min( (threadId % pairDistance)
const int leftId = min( (threadId % pairDistance)
+ (threadId / pairDistance) * blockWidth, count );
const int rightId = min( leftId + pairDistance, count );
@@ -147,7 +147,7 @@ __kernel
float2 greater = compareResult ? leftPt:rightPt;
float2 lesser = compareResult ? rightPt:leftPt;
corners[leftId] = sortOrder ? lesser : greater;
corners[rightId] = sortOrder ? greater : lesser;
}
@@ -195,20 +195,20 @@ __kernel
{
pt2 = scratch[j];
val2 = ELEM_FLT2(eig, pt2);
if(val2 > val1)
if(val2 > val1)
pos++;//calculate the rank of this element in this work group
else
else
{
if(val1 > val2)
continue;
else
else
{
// val1 and val2 are same
same++;
}
}
}
for (int j=0; j< same; j++)
for (int j=0; j< same; j++)
corners[pos + j] = pt1;
}
__kernel
@@ -240,15 +240,15 @@ __kernel
for(int k=0; k<wg; k++)
{
pt2 = corners[j*wg + k];
val2 = ELEM_FLT2(eig, pt2);
val2 = ELEM_FLT2(eig, pt2);
if(val1 > val2)
break;
else
{
//Increment only if the value is not the same.
//Increment only if the value is not the same.
if( val2 > val1 )
pos++;
else
else
same++;
}
}
@@ -257,20 +257,20 @@ __kernel
for(int k=0; k<remainder; k++)
{
pt2 = corners[(numOfGroups-1)*wg + k];
val2 = ELEM_FLT2(eig, pt2);
val2 = ELEM_FLT2(eig, pt2);
if(val1 > val2)
break;
else
{
//Don't increment if the value is the same.
//Don't increment if the value is the same.
//Two elements are same if (*userComp)(jData, iData) and (*userComp)(iData, jData) are both false
if(val2 > val1)
pos++;
else
else
same++;
}
}
for (int j=0; j< same; j++)
}
for (int j=0; j< same; j++)
corners[pos + j] = pt1;
}
+2 -2
View File
@@ -183,7 +183,7 @@ __kernel void warpAffineLinear_C1_D0(__global const uchar * restrict src, __glob
spos1 = src_offset + sy * srcStep + sx + 1;
spos2 = src_offset + (sy+1) * srcStep + sx;
spos3 = src_offset + (sy+1) * srcStep + sx + 1;
v0.s0 = scon0.s0 ? src[spos0.s0] : 0;
v1.s0 = scon1.s0 ? src[spos1.s0] : 0;
v2.s0 = scon2.s0 ? src[spos2.s0] : 0;
@@ -203,7 +203,7 @@ __kernel void warpAffineLinear_C1_D0(__global const uchar * restrict src, __glob
v1.s3 = scon1.s3 ? src[spos1.s3] : 0;
v2.s3 = scon2.s3 ? src[spos2.s3] : 0;
v3.s3 = scon3.s3 ? src[spos3.s3] : 0;
short4 itab0, itab1, itab2, itab3;
float4 taby, tabx;
taby = INTER_SCALE * convert_float4(ay);
@@ -116,7 +116,7 @@ __kernel void warpPerspectiveNN_C1_D0(__global uchar const * restrict src, __glo
sval.s1 = scon.s1 ? src[spos.s1] : 0;
sval.s2 = scon.s2 ? src[spos.s2] : 0;
sval.s3 = scon.s3 ? src[spos.s3] : 0;
dval = convert_uchar4(dcon) != (uchar4)(0,0,0,0) ? sval : dval;
dval = convert_uchar4(dcon) != (uchar4)(0,0,0,0) ? sval : dval;
*d = dval;
}
}
@@ -43,7 +43,7 @@
//
//M*/
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
#ifndef N // number of radices
#define N 4
@@ -71,14 +71,14 @@ __inline uint convertKey(uint converted_key)
converted_key ^= mask;
#elif defined(K_INT)
const uint SIGN_MASK = 1u << ((sizeof(int) * 8) - 1);
converted_key ^= SIGN_MASK;
converted_key ^= SIGN_MASK;
#else
#endif
return converted_key;
}
//FIXME(pengx17):
//FIXME(pengx17):
// exclusive scan, need to be optimized as this is too naive...
kernel
void naiveScanAddition(
@@ -108,7 +108,7 @@ kernel
{
const int RADIX_T = N;
const int RADICES_T = (1 << RADIX_T);
const int NUM_OF_ELEMENTS_PER_WORK_ITEM_T = RADICES_T;
const int NUM_OF_ELEMENTS_PER_WORK_ITEM_T = RADICES_T;
const int MASK_T = (1 << RADIX_T) - 1;
int localBuckets[16] = {0,0,0,0,0,0,0,0,
0,0,0,0,0,0,0,0};
+14 -14
View File
@@ -62,7 +62,7 @@
#endif
/////////////////////// Bitonic sort ////////////////////////////
// ported from
// ported from
// https://github.com/HSA-Libraries/Bolt/blob/master/include/bolt/cl/sort_by_key_kernels.cl
__kernel
void bitonicSort
@@ -82,7 +82,7 @@ __kernel
const int pairDistance = 1 << (stage - passOfStage);
const int blockWidth = 2 * pairDistance;
int leftId = min( (threadId % pairDistance)
int leftId = min( (threadId % pairDistance)
+ (threadId / pairDistance) * blockWidth, count );
int rightId = min( leftId + pairDistance, count );
@@ -90,7 +90,7 @@ __kernel
int temp;
const V_T lval = vals[leftId];
const V_T rval = vals[rightId];
const V_T rval = vals[rightId];
const K_T lkey = keys[leftId];
const K_T rkey = keys[rightId];
@@ -142,7 +142,7 @@ __kernel
int offset = groupID * wg;
int same = 0;
vals += offset;
keys += offset;
n = (groupID == (numOfGroups-1))? (count - wg*(numOfGroups-1)) : wg;
@@ -163,13 +163,13 @@ __kernel
for (int j=0;j<n;++j)
{
key2 = scratch[j];
if(my_comp(key2, key1))
if(my_comp(key2, key1))
pos++;//calculate the rank of this element in this work group
else
else
{
if(my_comp(key1, key2))
continue;
else
else
{
// key1 and key2 are same
same++;
@@ -209,15 +209,15 @@ __kernel
{
for(int k=0; k<wg; k++)
{
key2 = keys[j*wg + k];
key2 = keys[j*wg + k];
if(my_comp(key1, key2))
break;
else
{
//Increment only if the value is not the same.
//Increment only if the value is not the same.
if(my_comp(key2, key1))
pos++;
else
else
same++;
}
}
@@ -225,18 +225,18 @@ __kernel
for(int k=0; k<remainder; k++)
{
key2 = keys[(numOfGroups-1)*wg + k];
key2 = keys[(numOfGroups-1)*wg + k];
if(my_comp(key1, key2))
break;
else
{
//Don't increment if the value is the same.
//Don't increment if the value is the same.
if(my_comp(key2, key1))
pos++;
else
else
same++;
}
}
}
for (int j=0; j< same; j++)
{
vals[pos + j] = val1;
+42 -42
View File
@@ -56,7 +56,7 @@
// Use pre-computed gaussian and interp_weight lookup tables
__kernel void compute_hists_lut_kernel(
const int cblock_stride_x, const int cblock_stride_y,
const int cnbins, const int cblock_hist_size, const int img_block_width,
const int cnbins, const int cblock_hist_size, const int img_block_width,
const int blocks_in_group, const int blocks_total,
const int grad_quadstep, const int qangle_step,
__global const float* grad, __global const uchar* qangle,
@@ -76,20 +76,20 @@ __kernel void compute_hists_lut_kernel(
const int cell_y = lidY;
const int cell_thread_x = lidX - cell_x * 12;
__local float* hists = smem + lp * cnbins * (CELLS_PER_BLOCK_X *
__local float* hists = smem + lp * cnbins * (CELLS_PER_BLOCK_X *
CELLS_PER_BLOCK_Y * 12 + CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y);
__local float* final_hist = hists + cnbins *
__local float* final_hist = hists + cnbins *
(CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y * 12);
const int offset_x = gidX * cblock_stride_x + (cell_x << 2) + cell_thread_x;
const int offset_y = gidY * cblock_stride_y + (cell_y << 2);
__global const float* grad_ptr = (gid < blocks_total) ?
__global const float* grad_ptr = (gid < blocks_total) ?
grad + offset_y * grad_quadstep + (offset_x << 1) : grad;
__global const uchar* qangle_ptr = (gid < blocks_total) ?
qangle + offset_y * qangle_step + (offset_x << 1) : qangle;
__local float* hist = hists + 12 * (cell_y * CELLS_PER_BLOCK_Y + cell_x) +
__local float* hist = hists + 12 * (cell_y * CELLS_PER_BLOCK_Y + cell_x) +
cell_thread_x;
for (int bin_id = 0; bin_id < cnbins; ++bin_id)
hist[bin_id * 48] = 0.f;
@@ -130,7 +130,7 @@ __kernel void compute_hists_lut_kernel(
barrier(CLK_LOCAL_MEM_FENCE);
#endif
if (cell_thread_x == 0)
final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] =
final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] =
hist_[0] + hist_[1] + hist_[2];
}
#ifdef CPU
@@ -140,7 +140,7 @@ __kernel void compute_hists_lut_kernel(
int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 12 + cell_thread_x;
if ((tid < cblock_hist_size) && (gid < blocks_total))
{
__global float* block_hist = block_hists +
__global float* block_hist = block_hists +
(gidY * img_block_width + gidX) * cblock_hist_size;
block_hist[tid] = final_hist[tid];
}
@@ -149,7 +149,7 @@ __kernel void compute_hists_lut_kernel(
//-------------------------------------------------------------
// Normalization of histograms via L2Hys_norm
// optimized for the case of 9 bins
__kernel void normalize_hists_36_kernel(__global float* block_hists,
__kernel void normalize_hists_36_kernel(__global float* block_hists,
const float threshold, __local float *squares)
{
const int tid = get_local_id(0);
@@ -205,24 +205,24 @@ float reduce_smem(volatile __local float* smem, int size)
unsigned int tid = get_local_id(0);
float sum = smem[tid];
if (size >= 512) { if (tid < 256) smem[tid] = sum = sum + smem[tid + 256];
if (size >= 512) { if (tid < 256) smem[tid] = sum = sum + smem[tid + 256];
barrier(CLK_LOCAL_MEM_FENCE); }
if (size >= 256) { if (tid < 128) smem[tid] = sum = sum + smem[tid + 128];
if (size >= 256) { if (tid < 128) smem[tid] = sum = sum + smem[tid + 128];
barrier(CLK_LOCAL_MEM_FENCE); }
if (size >= 128) { if (tid < 64) smem[tid] = sum = sum + smem[tid + 64];
if (size >= 128) { if (tid < 64) smem[tid] = sum = sum + smem[tid + 64];
barrier(CLK_LOCAL_MEM_FENCE); }
#ifdef CPU
if (size >= 64) { if (tid < 32) smem[tid] = sum = sum + smem[tid + 32];
if (size >= 64) { if (tid < 32) smem[tid] = sum = sum + smem[tid + 32];
barrier(CLK_LOCAL_MEM_FENCE); }
if (size >= 32) { if (tid < 16) smem[tid] = sum = sum + smem[tid + 16];
barrier(CLK_LOCAL_MEM_FENCE); }
if (size >= 16) { if (tid < 8) smem[tid] = sum = sum + smem[tid + 8];
if (size >= 32) { if (tid < 16) smem[tid] = sum = sum + smem[tid + 16];
barrier(CLK_LOCAL_MEM_FENCE); }
if (size >= 8) { if (tid < 4) smem[tid] = sum = sum + smem[tid + 4];
if (size >= 16) { if (tid < 8) smem[tid] = sum = sum + smem[tid + 8];
barrier(CLK_LOCAL_MEM_FENCE); }
if (size >= 4) { if (tid < 2) smem[tid] = sum = sum + smem[tid + 2];
barrier(CLK_LOCAL_MEM_FENCE); }
if (size >= 2) { if (tid < 1) smem[tid] = sum = sum + smem[tid + 1];
if (size >= 8) { if (tid < 4) smem[tid] = sum = sum + smem[tid + 4];
barrier(CLK_LOCAL_MEM_FENCE); }
if (size >= 4) { if (tid < 2) smem[tid] = sum = sum + smem[tid + 2];
barrier(CLK_LOCAL_MEM_FENCE); }
if (size >= 2) { if (tid < 1) smem[tid] = sum = sum + smem[tid + 1];
barrier(CLK_LOCAL_MEM_FENCE); }
#else
if (tid < 32)
@@ -251,7 +251,7 @@ __kernel void normalize_hists_kernel(
const int gidX = get_group_id(0);
const int gidY = get_group_id(1);
__global float* hist = block_hists + (gidY * img_block_width + gidX) *
__global float* hist = block_hists + (gidY * img_block_width + gidX) *
block_hist_size + tid;
float elem = 0.f;
@@ -292,14 +292,14 @@ __kernel void classify_hists_180_kernel(
const int gidX = get_group_id(0);
const int gidY = get_group_id(1);
__global const float* hist = block_hists + (gidY * win_block_stride_y *
__global const float* hist = block_hists + (gidY * win_block_stride_y *
img_block_width + gidX * win_block_stride_x) * cblock_hist_size;
float product = 0.f;
for (int i = 0; i < cdescr_height; i++)
{
product += coefs[i * cdescr_width + tid] *
product += coefs[i * cdescr_width + tid] *
hist[i * img_block_width * cblock_hist_size + tid];
}
@@ -365,14 +365,14 @@ __kernel void classify_hists_252_kernel(
const int gidX = get_group_id(0);
const int gidY = get_group_id(1);
__global const float* hist = block_hists + (gidY * win_block_stride_y *
__global const float* hist = block_hists + (gidY * win_block_stride_y *
img_block_width + gidX * win_block_stride_x) * cblock_hist_size;
float product = 0.f;
if (tid < cdescr_width)
{
for (int i = 0; i < cdescr_height; i++)
product += coefs[i * cdescr_width + tid] *
product += coefs[i * cdescr_width + tid] *
hist[i * img_block_width * cblock_hist_size + tid];
}
@@ -402,7 +402,7 @@ __kernel void classify_hists_252_kernel(
barrier(CLK_LOCAL_MEM_FENCE);
#else
if (tid < 32)
{
{
smem[tid] = product = product + smem[tid + 32];
#if WAVE_SIZE < 32
} barrier(CLK_LOCAL_MEM_FENCE);
@@ -434,7 +434,7 @@ __kernel void classify_hists_kernel(
const int gidX = get_group_id(0);
const int gidY = get_group_id(1);
__global const float* hist = block_hists + (gidY * win_block_stride_y *
__global const float* hist = block_hists + (gidY * win_block_stride_y *
img_block_width + gidX * win_block_stride_x) * cblock_hist_size;
float product = 0.f;
@@ -442,7 +442,7 @@ __kernel void classify_hists_kernel(
{
int offset_y = i / cdescr_width;
int offset_x = i - offset_y * cdescr_width;
product += coefs[i] *
product += coefs[i] *
hist[offset_y * img_block_width * cblock_hist_size + offset_x];
}
@@ -472,7 +472,7 @@ __kernel void classify_hists_kernel(
barrier(CLK_LOCAL_MEM_FENCE);
#else
if (tid < 32)
{
{
smem[tid] = product = product + smem[tid + 32];
#if WAVE_SIZE < 32
} barrier(CLK_LOCAL_MEM_FENCE);
@@ -494,8 +494,8 @@ __kernel void classify_hists_kernel(
// Extract descriptors
__kernel void extract_descrs_by_rows_kernel(
const int cblock_hist_size, const int descriptors_quadstep,
const int cdescr_size, const int cdescr_width, const int img_block_width,
const int cblock_hist_size, const int descriptors_quadstep,
const int cdescr_size, const int cdescr_width, const int img_block_width,
const int win_block_stride_x, const int win_block_stride_y,
__global const float* block_hists, __global float* descriptors)
{
@@ -504,11 +504,11 @@ __kernel void extract_descrs_by_rows_kernel(
int gidY = get_group_id(1);
// Get left top corner of the window in src
__global const float* hist = block_hists + (gidY * win_block_stride_y *
__global const float* hist = block_hists + (gidY * win_block_stride_y *
img_block_width + gidX * win_block_stride_x) * cblock_hist_size;
// Get left top corner of the window in dst
__global float* descriptor = descriptors +
__global float* descriptor = descriptors +
(gidY * get_num_groups(0) + gidX) * descriptors_quadstep;
// Copy elements from src to dst
@@ -522,8 +522,8 @@ __kernel void extract_descrs_by_rows_kernel(
__kernel void extract_descrs_by_cols_kernel(
const int cblock_hist_size, const int descriptors_quadstep, const int cdescr_size,
const int cnblocks_win_x, const int cnblocks_win_y, const int img_block_width,
const int win_block_stride_x, const int win_block_stride_y,
const int cnblocks_win_x, const int cnblocks_win_y, const int img_block_width,
const int win_block_stride_x, const int win_block_stride_y,
__global const float* block_hists, __global float* descriptors)
{
int tid = get_local_id(0);
@@ -531,11 +531,11 @@ __kernel void extract_descrs_by_cols_kernel(
int gidY = get_group_id(1);
// Get left top corner of the window in src
__global const float* hist = block_hists + (gidY * win_block_stride_y *
__global const float* hist = block_hists + (gidY * win_block_stride_y *
img_block_width + gidX * win_block_stride_x) * cblock_hist_size;
// Get left top corner of the window in dst
__global float* descriptor = descriptors +
__global float* descriptor = descriptors +
(gidY * get_num_groups(0) + gidX) * descriptors_quadstep;
// Copy elements from src to dst
@@ -547,7 +547,7 @@ __kernel void extract_descrs_by_cols_kernel(
int y = block_idx / cnblocks_win_x;
int x = block_idx - y * cnblocks_win_x;
descriptor[(x * cnblocks_win_y + y) * cblock_hist_size + idx_in_block] =
descriptor[(x * cnblocks_win_y + y) * cblock_hist_size + idx_in_block] =
hist[(y * img_block_width + x) * cblock_hist_size + idx_in_block];
}
}
@@ -556,7 +556,7 @@ __kernel void extract_descrs_by_cols_kernel(
// Gradients computation
__kernel void compute_gradients_8UC4_kernel(
const int height, const int width,
const int height, const int width,
const int img_step, const int grad_quadstep, const int qangle_step,
const __global uchar4 * img, __global float * grad, __global uchar * qangle,
const float angle_scale, const char correct_gamma, const int cnbins)
@@ -600,9 +600,9 @@ __kernel void compute_gradients_8UC4_kernel(
barrier(CLK_LOCAL_MEM_FENCE);
if (x < width)
{
float3 a = (float3) (sh_row[tid], sh_row[tid + (NTHREADS + 2)],
float3 a = (float3) (sh_row[tid], sh_row[tid + (NTHREADS + 2)],
sh_row[tid + 2 * (NTHREADS + 2)]);
float3 b = (float3) (sh_row[tid + 2], sh_row[tid + 2 + (NTHREADS + 2)],
float3 b = (float3) (sh_row[tid + 2], sh_row[tid + 2 + (NTHREADS + 2)],
sh_row[tid + 2 + 2 * (NTHREADS + 2)]);
float3 dx;
@@ -659,7 +659,7 @@ __kernel void compute_gradients_8UC4_kernel(
}
__kernel void compute_gradients_8UC1_kernel(
const int height, const int width,
const int height, const int width,
const int img_step, const int grad_quadstep, const int qangle_step,
__global const uchar * img, __global float * grad, __global uchar * qangle,
const float angle_scale, const char correct_gamma, const int cnbins)
@@ -717,4 +717,4 @@ __kernel void compute_gradients_8UC1_kernel(
grad[ (gidY * grad_quadstep + x) << 1 ] = mag * (1.f - ang);
grad[ ((gidY * grad_quadstep + x) << 1) + 1 ] = mag * ang;
}
}
}
+1 -1
View File
@@ -190,7 +190,7 @@ __kernel void stereoKernel(__global unsigned char *left, __global unsigned char
{
int idx1 = y_tex * img_step + x_tex;
int idx2 = min(y_tex + ((radius << 1) + 1), cheight - 1) * img_step + x_tex;
barrier(CLK_LOCAL_MEM_FENCE);
StepDown(idx1, idx2, left, right, d, col_ssd);
+7 -7
View File
@@ -67,7 +67,7 @@
/////////////////////////////////////////////////////////////
T saturate_cast(float v){
#ifdef T_SHORT
return convert_short_sat_rte(v);
return convert_short_sat_rte(v);
#else
return v;
#endif
@@ -75,7 +75,7 @@ T saturate_cast(float v){
T4 saturate_cast4(float4 v){
#ifdef T_SHORT
return convert_short4_sat_rte(v);
return convert_short4_sat_rte(v);
#else
return v;
#endif
@@ -96,7 +96,7 @@ typedef struct
inline float pix_diff_1(const uchar4 l, __global const uchar *rs)
{
return abs((int)(l.x) - *rs);
return abs((int)(l.x) - *rs);
}
float pix_diff_4(const uchar4 l, __global const uchar *rs)
@@ -174,8 +174,8 @@ __kernel void comp_data(__global uchar *left, int left_rows, int left_cols, i
///////////////////////////////////////////////////////////////
//////////////////////// data step down ///////////////////////
///////////////////////////////////////////////////////////////
__kernel void data_step_down(__global T *src, int src_rows,
__global T *dst, int dst_rows, int dst_cols,
__kernel void data_step_down(__global T *src, int src_rows,
__global T *dst, int dst_rows, int dst_cols,
int src_step, int dst_step,
int cndisp)
{
@@ -289,7 +289,7 @@ void message(__global T *us_, __global T *ds_, __global T *ls_, __global T *rs_,
}
minimum += cmax_disc_term;
float4 sum = 0;
prev = convert_float4(t_dst[CNDISP - 1]);
for (int disp = CNDISP - 2; disp >= 0; disp--)
@@ -322,7 +322,7 @@ void message(__global T *us_, __global T *ds_, __global T *ls_, __global T *rs_,
__kernel void one_iteration(__global T *u, int u_step,
__global T *data, int data_step,
__global T *d, __global T *l, __global T *r,
int t, int cols, int rows,
int t, int cols, int rows,
float cmax_disc_term, float cdisc_single_jump)
{
const int y = get_global_id(1);

Some files were not shown because too many files have changed in this diff Show More