mirror of https://github.com/opencv/opencv.git
Conflicts: doc/tutorials/features2d/feature_detection/feature_detection.rst modules/bioinspired/doc/retina/index.rst modules/core/include/opencv2/core/core.hpp modules/core/include/opencv2/core/mat.hpp modules/core/include/opencv2/core/operations.hpp modules/core/src/stat.cpp modules/features2d/include/opencv2/features2d/features2d.hpp modules/imgproc/src/filter.cpp modules/legacy/src/dpstereo.cpp modules/nonfree/src/surf.ocl.cpp modules/ocl/doc/image_processing.rst modules/ocl/doc/object_detection.rst modules/ocl/include/opencv2/ocl/ocl.hpp modules/ocl/include/opencv2/ocl/private/util.hpp modules/ocl/src/arithm.cpp modules/ocl/src/canny.cpp modules/ocl/src/filtering.cpp modules/ocl/src/imgproc.cpp modules/ocl/src/initialization.cpp modules/ocl/src/matrix_operations.cpp modules/ocl/src/pyrdown.cpp modules/ocl/src/pyrup.cpp modules/ocl/src/split_merge.cpp modules/ocl/test/test_objdetect.cpp modules/ocl/test/utility.hpppull/1541/head
Roman Donchenko
12 years ago
commit
81eb4bd8be
142 changed files with 7748 additions and 13894 deletions
|
Before Width: | Height: | Size: 42 KiB |
|
Before Width: | Height: | Size: 12 KiB |
|
Before Width: | Height: | Size: 74 KiB |
|
Before Width: | Height: | Size: 42 KiB |
|
Before Width: | Height: | Size: 25 KiB |
|
Before Width: | Height: | Size: 25 KiB |
|
Before Width: | Height: | Size: 32 KiB |
|
Before Width: | Height: | Size: 14 KiB |
|
Before Width: | Height: | Size: 34 KiB |
|
Before Width: | Height: | Size: 38 KiB |
|
Before Width: | Height: | Size: 25 KiB |
|
After Width: | Height: | Size: 13 KiB |
|
After Width: | Height: | Size: 41 KiB |
|
After Width: | Height: | Size: 33 KiB |
|
After Width: | Height: | Size: 47 KiB |
|
After Width: | Height: | Size: 18 KiB |
|
After Width: | Height: | Size: 53 KiB |
|
After Width: | Height: | Size: 49 KiB |
|
After Width: | Height: | Size: 28 KiB |
|
After Width: | Height: | Size: 46 KiB |
|
After Width: | Height: | Size: 23 KiB |
|
After Width: | Height: | Size: 42 KiB |
|
After Width: | Height: | Size: 35 KiB |
@ -0,0 +1,110 @@ |
||||
|
||||
.. _Java_Eclipse: |
||||
|
||||
|
||||
Using OpenCV Java with Eclipse |
||||
********************************************* |
||||
|
||||
Since version 2.4.4 `OpenCV supports Java <http://opencv.org/opencv-java-api.html>`_. In this tutorial I will explain how to setup development environment for using OpenCV Java with Eclipse in **Windows**, so you can enjoy the benefits of garbage collected, very refactorable (rename variable, extract method and whatnot) modern language that enables you to write code with less effort and make less mistakes. Here we go. |
||||
|
||||
|
||||
Configuring Eclipse |
||||
=================== |
||||
|
||||
First, obtain a fresh release of OpenCV `from download page <http://opencv.org/downloads.html>`_ and extract it under a simple location like ``C:\OpenCV-2.4.6\``. I am using version 2.4.6, but the steps are more or less the same for other versions. |
||||
|
||||
Now, we will define OpenCV as a user library in Eclipse, so we can reuse the configuration for any project. Launch Eclipse and select :guilabel:`Window --> Preferences` from the menu. |
||||
|
||||
.. image:: images/1-window-preferences.png |
||||
:alt: Eclipse preferences |
||||
:align: center |
||||
|
||||
Navigate under :guilabel:`Java --> Build Path --> User Libraries` and click :guilabel:`New...`. |
||||
|
||||
.. image:: images/2-user-library-new.png |
||||
:alt: Creating a new library |
||||
:align: center |
||||
|
||||
Enter a name, e.g. ``OpenCV-2.4.6``, for your new library. |
||||
|
||||
.. image:: images/3-library-name.png |
||||
:alt: Naming the new library |
||||
:align: center |
||||
|
||||
Now select your new user library and click :guilabel:`Add External JARs...`. |
||||
|
||||
.. image:: images/4-add-external-jars.png |
||||
:alt: Adding external jar |
||||
:align: center |
||||
|
||||
Browse through ``C:\OpenCV-2.4.6\build\java\`` and select ``opencv-246.jar``. After adding the jar, extend the :guilabel:`opencv-246.jar` and select :guilabel:`Native library location` and press :guilabel:`Edit...`. |
||||
|
||||
.. image:: images/5-native-library.png |
||||
:alt: Selecting native library location 1 |
||||
:align: center |
||||
|
||||
Select :guilabel:`External Folder...` and browse to select the folder ``C:\OpenCV-2.4.6\build\java\x64``. If you have a 32-bit system you need to select the ``x86`` folder instead of ``x64``. |
||||
|
||||
.. image:: images/6-external-folder.png |
||||
:alt: Selecting native library location 2 |
||||
:align: center |
||||
|
||||
Your user library configuration should look like this: |
||||
|
||||
.. image:: images/7-user-library-final.png |
||||
:alt: Selecting native library location 2 |
||||
:align: center |
||||
|
||||
|
||||
Testing the configuration on a new Java project |
||||
===================================================== |
||||
|
||||
Now start creating a new Java project. |
||||
|
||||
.. image:: images/7_5-new-java-project.png |
||||
:alt: Creating new Java project |
||||
:align: center |
||||
|
||||
On the :guilabel:`Java Settings` step, under :guilabel:`Libraries` tab, select :guilabel:`Add Library...` and select :guilabel:`OpenCV-2.4.6`, then click :guilabel:`Finish`. |
||||
|
||||
.. image:: images/8-add-library.png |
||||
:alt: Adding user defined library 1 |
||||
:align: center |
||||
|
||||
.. image:: images/9-select-user-lib.png |
||||
:alt: Adding user defined library 2 |
||||
:align: center |
||||
|
||||
|
||||
Libraries should look like this: |
||||
|
||||
.. image:: images/10-new-project-created.png |
||||
:alt: Adding user defined library |
||||
:align: center |
||||
|
||||
|
||||
Now you have created and configured a new Java project it is time to test it. Create a new java file. Here is a starter code for your convenience: |
||||
|
||||
.. code-block:: java |
||||
|
||||
import org.opencv.core.Core; |
||||
import org.opencv.core.CvType; |
||||
import org.opencv.core.Mat; |
||||
|
||||
public class Hello |
||||
{ |
||||
public static void main( String[] args ) |
||||
{ |
||||
System.loadLibrary( Core.NATIVE_LIBRARY_NAME ); |
||||
Mat mat = Mat.eye( 3, 3, CvType.CV_8UC1 ); |
||||
System.out.println( "mat = " + mat.dump() ); |
||||
} |
||||
} |
||||
|
||||
When you run the code you should see 3x3 identity matrix as output. |
||||
|
||||
.. image:: images/11-the-code.png |
||||
:alt: Adding user defined library |
||||
:align: center |
||||
|
||||
That is it, whenever you start a new project just add the OpenCV user library that you have defined to your project and you are good to go. Enjoy your powerful, less painful development environment :) |
||||
|
After Width: | Height: | Size: 6.0 KiB |
@ -0,0 +1,334 @@ |
||||
Camera Calibration and 3D Reconstruction |
||||
======================================== |
||||
|
||||
.. highlight:: cpp |
||||
|
||||
|
||||
|
||||
ocl::StereoBM_OCL |
||||
--------------------- |
||||
.. ocv:class:: ocl::StereoBM_OCL |
||||
|
||||
Class computing stereo correspondence (disparity map) using the block matching algorithm. :: |
||||
|
||||
class CV_EXPORTS StereoBM_OCL |
||||
{ |
||||
public: |
||||
enum { BASIC_PRESET = 0, PREFILTER_XSOBEL = 1 }; |
||||
|
||||
enum { DEFAULT_NDISP = 64, DEFAULT_WINSZ = 19 }; |
||||
|
||||
//! the default constructor |
||||
StereoBM_OCL(); |
||||
//! the full constructor taking the camera-specific preset, number of disparities and the SAD window size. ndisparities must be multiple of 8. |
||||
StereoBM_OCL(int preset, int ndisparities = DEFAULT_NDISP, int winSize = DEFAULT_WINSZ); |
||||
|
||||
//! the stereo correspondence operator. Finds the disparity for the specified rectified stereo pair |
||||
//! Output disparity has CV_8U type. |
||||
void operator() ( const oclMat &left, const oclMat &right, oclMat &disparity); |
||||
|
||||
//! Some heuristics that tries to estmate |
||||
// if current GPU will be faster then CPU in this algorithm. |
||||
// It queries current active device. |
||||
static bool checkIfGpuCallReasonable(); |
||||
|
||||
int preset; |
||||
int ndisp; |
||||
int winSize; |
||||
|
||||
// If avergeTexThreshold == 0 => post procesing is disabled |
||||
// If avergeTexThreshold != 0 then disparity is set 0 in each point (x,y) where for left image |
||||
// SumOfHorizontalGradiensInWindow(x, y, winSize) < (winSize * winSize) * avergeTexThreshold |
||||
// i.e. input left image is low textured. |
||||
float avergeTexThreshold; |
||||
private: |
||||
/* hidden */ |
||||
}; |
||||
|
||||
|
||||
The class also performs pre- and post-filtering steps: Sobel pre-filtering (if ``PREFILTER_XSOBEL`` flag is set) and low textureness filtering (if ``averageTexThreshols > 0`` ). If ``avergeTexThreshold = 0`` , low textureness filtering is disabled. Otherwise, the disparity is set to 0 in each point ``(x, y)`` , where for the left image |
||||
|
||||
.. math:: |
||||
\sum HorizontalGradiensInWindow(x, y, winSize) < (winSize \cdot winSize) \cdot avergeTexThreshold |
||||
|
||||
This means that the input left image is low textured. |
||||
|
||||
|
||||
ocl::StereoBM_OCL::StereoBM_OCL |
||||
----------------------------------- |
||||
Enables :ocv:class:`ocl::StereoBM_OCL` constructors. |
||||
|
||||
.. ocv:function:: ocl::StereoBM_OCL::StereoBM_OCL() |
||||
|
||||
.. ocv:function:: ocl::StereoBM_OCL::StereoBM_OCL(int preset, int ndisparities = DEFAULT_NDISP, int winSize = DEFAULT_WINSZ) |
||||
|
||||
:param preset: Parameter presetting: |
||||
|
||||
* **BASIC_PRESET** Basic mode without pre-processing. |
||||
|
||||
* **PREFILTER_XSOBEL** Sobel pre-filtering mode. |
||||
|
||||
:param ndisparities: Number of disparities. It must be a multiple of 8 and less or equal to 256. |
||||
|
||||
:param winSize: Block size. |
||||
|
||||
|
||||
|
||||
ocl::StereoBM_OCL::operator () |
||||
---------------------------------- |
||||
Enables the stereo correspondence operator that finds the disparity for the specified rectified stereo pair. |
||||
|
||||
.. ocv:function:: void ocl::StereoBM_OCL::operator ()(const oclMat& left, const oclMat& right, oclMat& disparity) |
||||
|
||||
:param left: Left image. Only ``CV_8UC1`` type is supported. |
||||
|
||||
:param right: Right image with the same size and the same type as the left one. |
||||
|
||||
:param disparity: Output disparity map. It is a ``CV_8UC1`` image with the same size as the input images. |
||||
|
||||
:param stream: Stream for the asynchronous version. |
||||
|
||||
|
||||
ocl::StereoBM_OCL::checkIfGpuCallReasonable |
||||
----------------------------------------------- |
||||
Uses a heuristic method to estimate whether the current GPU is faster than the CPU in this algorithm. It queries the currently active device. |
||||
|
||||
.. ocv:function:: bool ocl::StereoBM_OCL::checkIfGpuCallReasonable() |
||||
|
||||
ocl::StereoBeliefPropagation |
||||
-------------------------------- |
||||
.. ocv:class:: ocl::StereoBeliefPropagation |
||||
|
||||
Class computing stereo correspondence using the belief propagation algorithm. :: |
||||
|
||||
class CV_EXPORTS StereoBeliefPropagation |
||||
{ |
||||
public: |
||||
enum { DEFAULT_NDISP = 64 }; |
||||
enum { DEFAULT_ITERS = 5 }; |
||||
enum { DEFAULT_LEVELS = 5 }; |
||||
static void estimateRecommendedParams(int width, int height, int &ndisp, int &iters, int &levels); |
||||
explicit StereoBeliefPropagation(int ndisp = DEFAULT_NDISP, |
||||
int iters = DEFAULT_ITERS, |
||||
int levels = DEFAULT_LEVELS, |
||||
int msg_type = CV_16S); |
||||
StereoBeliefPropagation(int ndisp, int iters, int levels, |
||||
float max_data_term, float data_weight, |
||||
float max_disc_term, float disc_single_jump, |
||||
int msg_type = CV_32F); |
||||
void operator()(const oclMat &left, const oclMat &right, oclMat &disparity); |
||||
void operator()(const oclMat &data, oclMat &disparity); |
||||
int ndisp; |
||||
int iters; |
||||
int levels; |
||||
float max_data_term; |
||||
float data_weight; |
||||
float max_disc_term; |
||||
float disc_single_jump; |
||||
int msg_type; |
||||
private: |
||||
/* hidden */ |
||||
}; |
||||
|
||||
The class implements algorithm described in [Felzenszwalb2006]_ . It can compute own data cost (using a truncated linear model) or use a user-provided data cost. |
||||
|
||||
.. note:: |
||||
|
||||
``StereoBeliefPropagation`` requires a lot of memory for message storage: |
||||
|
||||
.. math:: |
||||
|
||||
width \_ step \cdot height \cdot ndisp \cdot 4 \cdot (1 + 0.25) |
||||
|
||||
and for data cost storage: |
||||
|
||||
.. math:: |
||||
|
||||
width\_step \cdot height \cdot ndisp \cdot (1 + 0.25 + 0.0625 + \dotsm + \frac{1}{4^{levels}}) |
||||
|
||||
``width_step`` is the number of bytes in a line including padding. |
||||
|
||||
|
||||
|
||||
ocl::StereoBeliefPropagation::StereoBeliefPropagation |
||||
--------------------------------------------------------- |
||||
Enables the :ocv:class:`ocl::StereoBeliefPropagation` constructors. |
||||
|
||||
.. ocv:function:: ocl::StereoBeliefPropagation::StereoBeliefPropagation(int ndisp = DEFAULT_NDISP, int iters = DEFAULT_ITERS, int levels = DEFAULT_LEVELS, int msg_type = CV_16S) |
||||
|
||||
.. ocv:function:: ocl::StereoBeliefPropagation::StereoBeliefPropagation(int ndisp, int iters, int levels, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int msg_type = CV_32F) |
||||
|
||||
:param ndisp: Number of disparities. |
||||
|
||||
:param iters: Number of BP iterations on each level. |
||||
|
||||
:param levels: Number of levels. |
||||
|
||||
:param max_data_term: Threshold for data cost truncation. |
||||
|
||||
:param data_weight: Data weight. |
||||
|
||||
:param max_disc_term: Threshold for discontinuity truncation. |
||||
|
||||
:param disc_single_jump: Discontinuity single jump. |
||||
|
||||
:param msg_type: Type for messages. ``CV_16SC1`` and ``CV_32FC1`` types are supported. |
||||
|
||||
``StereoBeliefPropagation`` uses a truncated linear model for the data cost and discontinuity terms: |
||||
|
||||
.. math:: |
||||
|
||||
DataCost = data \_ weight \cdot \min ( \lvert Img_Left(x,y)-Img_Right(x-d,y) \rvert , max \_ data \_ term) |
||||
|
||||
.. math:: |
||||
|
||||
DiscTerm = \min (disc \_ single \_ jump \cdot \lvert f_1-f_2 \rvert , max \_ disc \_ term) |
||||
|
||||
For more details, see [Felzenszwalb2006]_. |
||||
|
||||
By default, :ocv:class:`ocl::StereoBeliefPropagation` uses floating-point arithmetics and the ``CV_32FC1`` type for messages. But it can also use fixed-point arithmetics and the ``CV_16SC1`` message type for better performance. To avoid an overflow in this case, the parameters must satisfy the following requirement: |
||||
|
||||
.. math:: |
||||
|
||||
10 \cdot 2^{levels-1} \cdot max \_ data \_ term < SHRT \_ MAX |
||||
|
||||
|
||||
|
||||
ocl::StereoBeliefPropagation::estimateRecommendedParams |
||||
----------------------------------------------------------- |
||||
Uses a heuristic method to compute the recommended parameters ( ``ndisp``, ``iters`` and ``levels`` ) for the specified image size ( ``width`` and ``height`` ). |
||||
|
||||
.. ocv:function:: void ocl::StereoBeliefPropagation::estimateRecommendedParams(int width, int height, int& ndisp, int& iters, int& levels) |
||||
|
||||
|
||||
|
||||
ocl::StereoBeliefPropagation::operator () |
||||
--------------------------------------------- |
||||
Enables the stereo correspondence operator that finds the disparity for the specified rectified stereo pair or data cost. |
||||
|
||||
.. ocv:function:: void ocl::StereoBeliefPropagation::operator ()(const oclMat& left, const oclMat& right, oclMat& disparity) |
||||
|
||||
.. ocv:function:: void ocl::StereoBeliefPropagation::operator ()(const oclMat& data, oclMat& disparity) |
||||
|
||||
:param left: Left image. ``CV_8UC1`` , ``CV_8UC3`` and ``CV_8UC4`` types are supported. |
||||
|
||||
:param right: Right image with the same size and the same type as the left one. |
||||
|
||||
:param data: User-specified data cost, a matrix of ``msg_type`` type and ``Size(<image columns>*ndisp, <image rows>)`` size. |
||||
|
||||
:param disparity: Output disparity map. If ``disparity`` is empty, the output type is ``CV_16SC1`` . Otherwise, the type is retained. |
||||
|
||||
:param stream: Stream for the asynchronous version. |
||||
|
||||
ocl::StereoConstantSpaceBP |
||||
------------------------------ |
||||
.. ocv:class:: ocl::StereoConstantSpaceBP |
||||
|
||||
Class computing stereo correspondence using the constant space belief propagation algorithm. :: |
||||
|
||||
class CV_EXPORTS StereoConstantSpaceBP |
||||
{ |
||||
public: |
||||
enum { DEFAULT_NDISP = 128 }; |
||||
enum { DEFAULT_ITERS = 8 }; |
||||
enum { DEFAULT_LEVELS = 4 }; |
||||
enum { DEFAULT_NR_PLANE = 4 }; |
||||
static void estimateRecommendedParams(int width, int height, int &ndisp, int &iters, int &levels, int &nr_plane); |
||||
explicit StereoConstantSpaceBP( |
||||
int ndisp = DEFAULT_NDISP, |
||||
int iters = DEFAULT_ITERS, |
||||
int levels = DEFAULT_LEVELS, |
||||
int nr_plane = DEFAULT_NR_PLANE, |
||||
int msg_type = CV_32F); |
||||
StereoConstantSpaceBP(int ndisp, int iters, int levels, int nr_plane, |
||||
float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, |
||||
int min_disp_th = 0, |
||||
int msg_type = CV_32F); |
||||
void operator()(const oclMat &left, const oclMat &right, oclMat &disparity); |
||||
int ndisp; |
||||
int iters; |
||||
int levels; |
||||
int nr_plane; |
||||
float max_data_term; |
||||
float data_weight; |
||||
float max_disc_term; |
||||
float disc_single_jump; |
||||
int min_disp_th; |
||||
int msg_type; |
||||
bool use_local_init_data_cost; |
||||
private: |
||||
/* hidden */ |
||||
}; |
||||
|
||||
The class implements algorithm described in [Yang2010]_. ``StereoConstantSpaceBP`` supports both local minimum and global minimum data cost initialization algorithms. For more details, see the paper mentioned above. By default, a local algorithm is used. To enable a global algorithm, set ``use_local_init_data_cost`` to ``false`` . |
||||
|
||||
|
||||
ocl::StereoConstantSpaceBP::StereoConstantSpaceBP |
||||
----------------------------------------------------- |
||||
Enables the :ocv:class:`ocl::StereoConstantSpaceBP` constructors. |
||||
|
||||
.. ocv:function:: ocl::StereoConstantSpaceBP::StereoConstantSpaceBP(int ndisp = DEFAULT_NDISP, int iters = DEFAULT_ITERS, int levels = DEFAULT_LEVELS, int nr_plane = DEFAULT_NR_PLANE, int msg_type = CV_32F) |
||||
|
||||
.. ocv:function:: ocl::StereoConstantSpaceBP::StereoConstantSpaceBP(int ndisp, int iters, int levels, int nr_plane, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th = 0, int msg_type = CV_32F) |
||||
|
||||
:param ndisp: Number of disparities. |
||||
|
||||
:param iters: Number of BP iterations on each level. |
||||
|
||||
:param levels: Number of levels. |
||||
|
||||
:param nr_plane: Number of disparity levels on the first level. |
||||
|
||||
:param max_data_term: Truncation of data cost. |
||||
|
||||
:param data_weight: Data weight. |
||||
|
||||
:param max_disc_term: Truncation of discontinuity. |
||||
|
||||
:param disc_single_jump: Discontinuity single jump. |
||||
|
||||
:param min_disp_th: Minimal disparity threshold. |
||||
|
||||
:param msg_type: Type for messages. ``CV_16SC1`` and ``CV_32FC1`` types are supported. |
||||
|
||||
``StereoConstantSpaceBP`` uses a truncated linear model for the data cost and discontinuity terms: |
||||
|
||||
.. math:: |
||||
|
||||
DataCost = data \_ weight \cdot \min ( \lvert I_2-I_1 \rvert , max \_ data \_ term) |
||||
|
||||
.. math:: |
||||
|
||||
DiscTerm = \min (disc \_ single \_ jump \cdot \lvert f_1-f_2 \rvert , max \_ disc \_ term) |
||||
|
||||
For more details, see [Yang2010]_. |
||||
|
||||
By default, ``StereoConstantSpaceBP`` uses floating-point arithmetics and the ``CV_32FC1`` type for messages. But it can also use fixed-point arithmetics and the ``CV_16SC1`` message type for better performance. To avoid an overflow in this case, the parameters must satisfy the following requirement: |
||||
|
||||
.. math:: |
||||
|
||||
10 \cdot 2^{levels-1} \cdot max \_ data \_ term < SHRT \_ MAX |
||||
|
||||
|
||||
|
||||
ocl::StereoConstantSpaceBP::estimateRecommendedParams |
||||
--------------------------------------------------------- |
||||
Uses a heuristic method to compute parameters (ndisp, iters, levelsand nrplane) for the specified image size (widthand height). |
||||
|
||||
.. ocv:function:: void ocl::StereoConstantSpaceBP::estimateRecommendedParams(int width, int height, int& ndisp, int& iters, int& levels, int& nr_plane) |
||||
|
||||
|
||||
|
||||
ocl::StereoConstantSpaceBP::operator () |
||||
------------------------------------------- |
||||
Enables the stereo correspondence operator that finds the disparity for the specified rectified stereo pair. |
||||
|
||||
.. ocv:function:: void ocl::StereoConstantSpaceBP::operator ()(const oclMat& left, const oclMat& right, oclMat& disparity) |
||||
|
||||
:param left: Left image. ``CV_8UC1`` , ``CV_8UC3`` and ``CV_8UC4`` types are supported. |
||||
|
||||
:param right: Right image with the same size and the same type as the left one. |
||||
|
||||
:param disparity: Output disparity map. If ``disparity`` is empty, the output type is ``CV_16SC1`` . Otherwise, the output type is ``disparity.type()`` . |
||||
|
||||
:param stream: Stream for the asynchronous version. |
||||
|
After Width: | Height: | Size: 64 KiB |
@ -0,0 +1,88 @@ |
||||
ml.Machine Learning |
||||
============================= |
||||
|
||||
.. highlight:: cpp |
||||
|
||||
ocl::KNearestNeighbour |
||||
-------------------------- |
||||
.. ocv:class:: ocl::KNearestNeighbour : public ocl::CvKNearest |
||||
|
||||
The class implements K-Nearest Neighbors model as described in the beginning of this section. |
||||
|
||||
ocl::KNearestNeighbour |
||||
-------------------------- |
||||
Computes the weighted sum of two arrays. :: |
||||
|
||||
class CV_EXPORTS KNearestNeighbour: public CvKNearest |
||||
{ |
||||
public: |
||||
KNearestNeighbour(); |
||||
~KNearestNeighbour(); |
||||
|
||||
bool train(const Mat& trainData, Mat& labels, Mat& sampleIdx = Mat().setTo(Scalar::all(0)), |
||||
bool isRegression = false, int max_k = 32, bool updateBase = false); |
||||
|
||||
void clear(); |
||||
|
||||
void find_nearest(const oclMat& samples, int k, oclMat& lables); |
||||
|
||||
private: |
||||
/* hidden */ |
||||
}; |
||||
|
||||
ocl::KNearestNeighbour::train |
||||
--------------------------------- |
||||
Trains the model. |
||||
|
||||
.. ocv:function:: bool ocl::KNearestNeighbour::train(const Mat& trainData, Mat& labels, Mat& sampleIdx = Mat().setTo(Scalar::all(0)), bool isRegression = false, int max_k = 32, bool updateBase = false) |
||||
|
||||
:param isRegression: Type of the problem: ``true`` for regression and ``false`` for classification. |
||||
|
||||
:param maxK: Number of maximum neighbors that may be passed to the method :ocv:func:`CvKNearest::find_nearest`. |
||||
|
||||
: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: |
||||
|
||||
* Only ``CV_ROW_SAMPLE`` data layout is supported. |
||||
* Input variables are all ordered. |
||||
* Output variables can be either categorical ( ``is_regression=false`` ) or ordered ( ``is_regression=true`` ). |
||||
* Variable subsets (``var_idx``) and missing measurements are not supported. |
||||
|
||||
ocl::KNearestNeighbour::find_nearest |
||||
---------------------------------------- |
||||
Finds the neighbors and predicts responses for input vectors. |
||||
|
||||
.. ocv:function:: void ocl::KNearestNeighbour::find_nearest(const oclMat& samples, int k, oclMat& lables ) |
||||
|
||||
:param samples: Input samples stored by rows. It is a single-precision floating-point matrix of :math:`number\_of\_samples \times number\_of\_features` size. |
||||
|
||||
:param k: Number of used nearest neighbors. It must satisfy constraint: :math:`k \le` :ocv:func:`CvKNearest::get_max_k`. |
||||
|
||||
:param labels: Vector with results of prediction (regression or classification) for each input sample. It is a single-precision floating-point vector with ``number_of_samples`` elements. |
||||
|
||||
ocl::kmeans |
||||
--------------- |
||||
Finds centers of clusters and groups input samples around the clusters. |
||||
|
||||
.. ocv:function:: double ocl::kmeans(const oclMat &src, int K, oclMat &bestLabels, TermCriteria criteria, int attemps, int flags, oclMat ¢ers) |
||||
|
||||
:param src: Floating-point matrix of input samples, one row per sample. |
||||
|
||||
:param K: Number of clusters to split the set by. |
||||
|
||||
:param bestLabels: Input/output integer array that stores the cluster indices for every sample. |
||||
|
||||
:param criteria: The algorithm termination criteria, that is, the maximum number of iterations and/or the desired accuracy. The accuracy is specified as ``criteria.epsilon``. As soon as each of the cluster centers moves by less than ``criteria.epsilon`` on some iteration, the algorithm stops. |
||||
|
||||
:param attempts: Flag to specify the number of times the algorithm is executed using different initial labellings. The algorithm returns the labels that yield the best compactness (see the last function parameter). |
||||
|
||||
:param flags: Flag that can take the following values: |
||||
|
||||
* **KMEANS_RANDOM_CENTERS** Select random initial centers in each attempt. |
||||
|
||||
* **KMEANS_PP_CENTERS** Use ``kmeans++`` center initialization by Arthur and Vassilvitskii [Arthur2007]. |
||||
|
||||
* **KMEANS_USE_INITIAL_LABELS** During the first (and possibly the only) attempt, use the user-supplied labels instead of computing them from the initial centers. For the second and further attempts, use the random or semi-random centers. Use one of ``KMEANS_*_CENTERS`` flag to specify the exact method. |
||||
|
||||
:param centers: Output matrix of the cluster centers, one row per each cluster center. |
||||
@ -0,0 +1,570 @@ |
||||
Video Analysis |
||||
============================= |
||||
|
||||
.. highlight:: cpp |
||||
|
||||
ocl::GoodFeaturesToTrackDetector_OCL |
||||
---------------------------------------- |
||||
.. ocv:class:: ocl::GoodFeaturesToTrackDetector_OCL |
||||
|
||||
Class used for strong corners detection on an image. :: |
||||
|
||||
class GoodFeaturesToTrackDetector_OCL |
||||
{ |
||||
public: |
||||
explicit GoodFeaturesToTrackDetector_OCL(int maxCorners = 1000, double qualityLevel = 0.01, double minDistance = 0.0, |
||||
int blockSize = 3, bool useHarrisDetector = false, double harrisK = 0.04); |
||||
|
||||
//! return 1 rows matrix with CV_32FC2 type |
||||
void operator ()(const oclMat& image, oclMat& corners, const oclMat& mask = oclMat()); |
||||
//! download points of type Point2f to a vector. the vector's content will be erased |
||||
void downloadPoints(const oclMat &points, std::vector<Point2f> &points_v); |
||||
|
||||
int maxCorners; |
||||
double qualityLevel; |
||||
double minDistance; |
||||
|
||||
int blockSize; |
||||
bool useHarrisDetector; |
||||
double harrisK; |
||||
void releaseMemory() |
||||
{ |
||||
Dx_.release(); |
||||
Dy_.release(); |
||||
eig_.release(); |
||||
minMaxbuf_.release(); |
||||
tmpCorners_.release(); |
||||
} |
||||
}; |
||||
|
||||
The class finds the most prominent corners in the image. |
||||
|
||||
.. seealso:: :ocv:func:`goodFeaturesToTrack()` |
||||
|
||||
ocl::GoodFeaturesToTrackDetector_OCL::GoodFeaturesToTrackDetector_OCL |
||||
------------------------------------------------------------------------- |
||||
Constructor. |
||||
|
||||
.. ocv:function:: ocl::GoodFeaturesToTrackDetector_OCL::GoodFeaturesToTrackDetector_OCL(int maxCorners = 1000, double qualityLevel = 0.01, double minDistance = 0.0, int blockSize = 3, bool useHarrisDetector = false, double harrisK = 0.04) |
||||
|
||||
:param maxCorners: Maximum number of corners to return. If there are more corners than are found, the strongest of them is returned. |
||||
|
||||
:param qualityLevel: Parameter characterizing the minimal accepted quality of image corners. The parameter value is multiplied by the best corner quality measure, which is the minimal eigenvalue (see :ocv:func:`ocl::cornerMinEigenVal` ) or the Harris function response (see :ocv:func:`ocl::cornerHarris` ). The corners with the quality measure less than the product are rejected. For example, if the best corner has the quality measure = 1500, and the ``qualityLevel=0.01`` , then all the corners with the quality measure less than 15 are rejected. |
||||
|
||||
:param minDistance: Minimum possible Euclidean distance between the returned corners. |
||||
|
||||
:param blockSize: Size of an average block for computing a derivative covariation matrix over each pixel neighborhood. See :ocv:func:`cornerEigenValsAndVecs` . |
||||
|
||||
:param useHarrisDetector: Parameter indicating whether to use a Harris detector (see :ocv:func:`ocl::cornerHarris`) or :ocv:func:`ocl::cornerMinEigenVal`. |
||||
|
||||
:param harrisK: Free parameter of the Harris detector. |
||||
|
||||
ocl::GoodFeaturesToTrackDetector_OCL::operator () |
||||
------------------------------------------------------- |
||||
Finds the most prominent corners in the image. |
||||
|
||||
.. ocv:function:: void ocl::GoodFeaturesToTrackDetector_OCL::operator ()(const oclMat& image, oclMat& corners, const oclMat& mask = oclMat()) |
||||
|
||||
:param image: Input 8-bit, single-channel image. |
||||
|
||||
:param corners: Output vector of detected corners (it will be one row matrix with CV_32FC2 type). |
||||
|
||||
:param mask: Optional region of interest. If the image is not empty (it needs to have the type ``CV_8UC1`` and the same size as ``image`` ), it specifies the region in which the corners are detected. |
||||
|
||||
.. seealso:: :ocv:func:`goodFeaturesToTrack` |
||||
|
||||
ocl::GoodFeaturesToTrackDetector_OCL::releaseMemory |
||||
-------------------------------------------------------- |
||||
Releases inner buffers memory. |
||||
|
||||
.. ocv:function:: void ocl::GoodFeaturesToTrackDetector_OCL::releaseMemory() |
||||
|
||||
ocl::FarnebackOpticalFlow |
||||
------------------------------- |
||||
.. ocv:class:: ocl::FarnebackOpticalFlow |
||||
|
||||
Class computing a dense optical flow using the Gunnar Farneback's algorithm. :: |
||||
|
||||
class CV_EXPORTS FarnebackOpticalFlow |
||||
{ |
||||
public: |
||||
FarnebackOpticalFlow(); |
||||
|
||||
int numLevels; |
||||
double pyrScale; |
||||
bool fastPyramids; |
||||
int winSize; |
||||
int numIters; |
||||
int polyN; |
||||
double polySigma; |
||||
int flags; |
||||
|
||||
void operator ()(const oclMat &frame0, const oclMat &frame1, oclMat &flowx, oclMat &flowy); |
||||
|
||||
void releaseMemory(); |
||||
|
||||
private: |
||||
/* hidden */ |
||||
}; |
||||
|
||||
ocl::FarnebackOpticalFlow::operator () |
||||
------------------------------------------ |
||||
Computes a dense optical flow using the Gunnar Farneback's algorithm. |
||||
|
||||
.. ocv:function:: void ocl::FarnebackOpticalFlow::operator ()(const oclMat &frame0, const oclMat &frame1, oclMat &flowx, oclMat &flowy) |
||||
|
||||
:param frame0: First 8-bit gray-scale input image |
||||
:param frame1: Second 8-bit gray-scale input image |
||||
:param flowx: Flow horizontal component |
||||
:param flowy: Flow vertical component |
||||
:param s: Stream |
||||
|
||||
.. seealso:: :ocv:func:`calcOpticalFlowFarneback` |
||||
|
||||
ocl::FarnebackOpticalFlow::releaseMemory |
||||
-------------------------------------------- |
||||
Releases unused auxiliary memory buffers. |
||||
|
||||
.. ocv:function:: void ocl::FarnebackOpticalFlow::releaseMemory() |
||||
|
||||
|
||||
ocl::PyrLKOpticalFlow |
||||
------------------------- |
||||
.. ocv:class:: ocl::PyrLKOpticalFlow |
||||
|
||||
Class used for calculating an optical flow. :: |
||||
|
||||
class PyrLKOpticalFlow |
||||
{ |
||||
public: |
||||
PyrLKOpticalFlow(); |
||||
|
||||
void sparse(const oclMat& prevImg, const oclMat& nextImg, const oclMat& prevPts, oclMat& nextPts, |
||||
oclMat& status, oclMat* err = 0); |
||||
|
||||
void dense(const oclMat& prevImg, const oclMat& nextImg, oclMat& u, oclMat& v, oclMat* err = 0); |
||||
|
||||
Size winSize; |
||||
int maxLevel; |
||||
int iters; |
||||
double derivLambda; |
||||
bool useInitialFlow; |
||||
float minEigThreshold; |
||||
bool getMinEigenVals; |
||||
|
||||
void releaseMemory(); |
||||
|
||||
private: |
||||
/* hidden */ |
||||
}; |
||||
|
||||
The class can calculate an optical flow for a sparse feature set or dense optical flow using the iterative Lucas-Kanade method with pyramids. |
||||
|
||||
.. seealso:: :ocv:func:`calcOpticalFlowPyrLK` |
||||
|
||||
ocl::PyrLKOpticalFlow::sparse |
||||
--------------------------------- |
||||
Calculate an optical flow for a sparse feature set. |
||||
|
||||
.. ocv:function:: void ocl::PyrLKOpticalFlow::sparse(const oclMat& prevImg, const oclMat& nextImg, const oclMat& prevPts, oclMat& nextPts, oclMat& status, oclMat* err = 0) |
||||
|
||||
:param prevImg: First 8-bit input image (supports both grayscale and color images). |
||||
|
||||
:param nextImg: Second input image of the same size and the same type as ``prevImg`` . |
||||
|
||||
:param prevPts: Vector of 2D points for which the flow needs to be found. It must be one row matrix with CV_32FC2 type. |
||||
|
||||
:param nextPts: Output vector of 2D points (with single-precision floating-point coordinates) containing the calculated new positions of input features in the second image. When ``useInitialFlow`` is true, the vector must have the same size as in the input. |
||||
|
||||
:param status: Output status vector (CV_8UC1 type). Each element of the vector is set to 1 if the flow for the corresponding features has been found. Otherwise, it is set to 0. |
||||
|
||||
:param err: Output vector (CV_32FC1 type) that contains the difference between patches around the original and moved points or min eigen value if ``getMinEigenVals`` is checked. It can be NULL, if not needed. |
||||
|
||||
.. seealso:: :ocv:func:`calcOpticalFlowPyrLK` |
||||
|
||||
|
||||
ocl::PyrLKOpticalFlow::dense |
||||
--------------------------------- |
||||
Calculate dense optical flow. |
||||
|
||||
.. ocv:function:: void ocl::PyrLKOpticalFlow::dense(const oclMat& prevImg, const oclMat& nextImg, oclMat& u, oclMat& v, oclMat* err = 0) |
||||
|
||||
:param prevImg: First 8-bit grayscale input image. |
||||
|
||||
:param nextImg: Second input image of the same size and the same type as ``prevImg`` . |
||||
|
||||
:param u: Horizontal component of the optical flow of the same size as input images, 32-bit floating-point, single-channel |
||||
|
||||
:param v: Vertical component of the optical flow of the same size as input images, 32-bit floating-point, single-channel |
||||
|
||||
:param err: Output vector (CV_32FC1 type) that contains the difference between patches around the original and moved points or min eigen value if ``getMinEigenVals`` is checked. It can be NULL, if not needed. |
||||
|
||||
|
||||
ocl::PyrLKOpticalFlow::releaseMemory |
||||
---------------------------------------- |
||||
Releases inner buffers memory. |
||||
|
||||
.. ocv:function:: void ocl::PyrLKOpticalFlow::releaseMemory() |
||||
|
||||
ocl::interpolateFrames |
||||
-------------------------- |
||||
Interpolates frames (images) using provided optical flow (displacement field). |
||||
|
||||
.. ocv:function:: void ocl::interpolateFrames(const oclMat& frame0, const oclMat& frame1, const oclMat& fu, const oclMat& fv, const oclMat& bu, const oclMat& bv, float pos, oclMat& newFrame, oclMat& buf) |
||||
|
||||
:param frame0: First frame (32-bit floating point images, single channel). |
||||
|
||||
:param frame1: Second frame. Must have the same type and size as ``frame0`` . |
||||
|
||||
:param fu: Forward horizontal displacement. |
||||
|
||||
:param fv: Forward vertical displacement. |
||||
|
||||
:param bu: Backward horizontal displacement. |
||||
|
||||
:param bv: Backward vertical displacement. |
||||
|
||||
:param pos: New frame position. |
||||
|
||||
:param newFrame: Output image. |
||||
|
||||
:param buf: Temporary buffer, will have width x 6*height size, CV_32FC1 type and contain 6 oclMat: occlusion masks for first frame, occlusion masks for second, interpolated forward horizontal flow, interpolated forward vertical flow, interpolated backward horizontal flow, interpolated backward vertical flow. |
||||
|
||||
:param stream: Stream for the asynchronous version. |
||||
|
||||
ocl::KalmanFilter |
||||
-------------------- |
||||
.. ocv:class:: ocl::KalmanFilter |
||||
|
||||
Kalman filter class. :: |
||||
|
||||
class CV_EXPORTS KalmanFilter |
||||
{ |
||||
public: |
||||
KalmanFilter(); |
||||
//! the full constructor taking the dimensionality of the state, of the measurement and of the control vector |
||||
KalmanFilter(int dynamParams, int measureParams, int controlParams=0, int type=CV_32F); |
||||
//! re-initializes Kalman filter. The previous content is destroyed. |
||||
void init(int dynamParams, int measureParams, int controlParams=0, int type=CV_32F); |
||||
|
||||
const oclMat& predict(const oclMat& control=oclMat()); |
||||
const oclMat& correct(const oclMat& measurement); |
||||
|
||||
oclMat statePre; //!< predicted state (x'(k)): x(k)=A*x(k-1)+B*u(k) |
||||
oclMat statePost; //!< corrected state (x(k)): x(k)=x'(k)+K(k)*(z(k)-H*x'(k)) |
||||
oclMat transitionMatrix; //!< state transition matrix (A) |
||||
oclMat controlMatrix; //!< control matrix (B) (not used if there is no control) |
||||
oclMat measurementMatrix; //!< measurement matrix (H) |
||||
oclMat processNoiseCov; //!< process noise covariance matrix (Q) |
||||
oclMat measurementNoiseCov;//!< measurement noise covariance matrix (R) |
||||
oclMat errorCovPre; //!< priori error estimate covariance matrix (P'(k)): P'(k)=A*P(k-1)*At + Q)*/ |
||||
oclMat gain; //!< Kalman gain matrix (K(k)): K(k)=P'(k)*Ht*inv(H*P'(k)*Ht+R) |
||||
oclMat errorCovPost; //!< posteriori error estimate covariance matrix (P(k)): P(k)=(I-K(k)*H)*P'(k) |
||||
private: |
||||
/* hidden */ |
||||
}; |
||||
|
||||
ocl::KalmanFilter::KalmanFilter |
||||
---------------------------------- |
||||
The constructors. |
||||
|
||||
.. ocv:function:: ocl::KalmanFilter::KalmanFilter() |
||||
|
||||
.. ocv:function:: ocl::KalmanFilter::KalmanFilter(int dynamParams, int measureParams, int controlParams=0, int type=CV_32F) |
||||
|
||||
The full constructor. |
||||
|
||||
:param dynamParams: Dimensionality of the state. |
||||
|
||||
:param measureParams: Dimensionality of the measurement. |
||||
|
||||
:param controlParams: Dimensionality of the control vector. |
||||
|
||||
:param type: Type of the created matrices that should be ``CV_32F`` or ``CV_64F``. |
||||
|
||||
|
||||
ocl::KalmanFilter::init |
||||
--------------------------- |
||||
Re-initializes Kalman filter. The previous content is destroyed. |
||||
|
||||
.. ocv:function:: void ocl::KalmanFilter::init(int dynamParams, int measureParams, int controlParams=0, int type=CV_32F) |
||||
|
||||
:param dynamParams: Dimensionalityensionality of the state. |
||||
|
||||
:param measureParams: Dimensionality of the measurement. |
||||
|
||||
:param controlParams: Dimensionality of the control vector. |
||||
|
||||
:param type: Type of the created matrices that should be ``CV_32F`` or ``CV_64F``. |
||||
|
||||
|
||||
ocl::KalmanFilter::predict |
||||
------------------------------ |
||||
Computes a predicted state. |
||||
|
||||
.. ocv:function:: const oclMat& ocl::KalmanFilter::predict(const oclMat& control=oclMat()) |
||||
|
||||
:param control: The optional input control |
||||
|
||||
|
||||
ocl::KalmanFilter::correct |
||||
----------------------------- |
||||
Updates the predicted state from the measurement. |
||||
|
||||
.. ocv:function:: const oclMat& ocl::KalmanFilter::correct(const oclMat& measurement) |
||||
|
||||
:param measurement: The measured system parameters |
||||
|
||||
|
||||
ocl::BackgroundSubtractor |
||||
---------------------------- |
||||
.. ocv:class:: ocl::BackgroundSubtractor |
||||
|
||||
Base class for background/foreground segmentation. :: |
||||
|
||||
class CV_EXPORTS BackgroundSubtractor |
||||
{ |
||||
public: |
||||
//! the virtual destructor |
||||
virtual ~BackgroundSubtractor(); |
||||
//! the update operator that takes the next video frame and returns the current foreground mask as 8-bit binary image. |
||||
virtual void operator()(const oclMat& image, oclMat& fgmask, float learningRate); |
||||
|
||||
//! computes a background image |
||||
virtual void getBackgroundImage(oclMat& backgroundImage) const = 0; |
||||
}; |
||||
|
||||
|
||||
The class is only used to define the common interface for the whole family of background/foreground segmentation algorithms. |
||||
|
||||
|
||||
ocl::BackgroundSubtractor::operator() |
||||
----------------------------------------- |
||||
Computes a foreground mask. |
||||
|
||||
.. ocv:function:: void ocl::BackgroundSubtractor::operator()(const oclMat& image, oclMat& fgmask, float learningRate) |
||||
|
||||
:param image: Next video frame. |
||||
|
||||
:param fgmask: The output foreground mask as an 8-bit binary image. |
||||
|
||||
|
||||
ocl::BackgroundSubtractor::getBackgroundImage |
||||
------------------------------------------------- |
||||
Computes a background image. |
||||
|
||||
.. ocv:function:: void ocl::BackgroundSubtractor::getBackgroundImage(oclMat& backgroundImage) const |
||||
|
||||
:param backgroundImage: The output background image. |
||||
|
||||
.. note:: Sometimes the background image can be very blurry, as it contain the average background statistics. |
||||
|
||||
ocl::MOG |
||||
------------ |
||||
.. ocv:class:: ocl::MOG : public ocl::BackgroundSubtractor |
||||
|
||||
Gaussian Mixture-based Backbround/Foreground Segmentation Algorithm. :: |
||||
|
||||
class CV_EXPORTS MOG: public cv::ocl::BackgroundSubtractor |
||||
{ |
||||
public: |
||||
//! the default constructor |
||||
MOG(int nmixtures = -1); |
||||
|
||||
//! re-initiaization method |
||||
void initialize(Size frameSize, int frameType); |
||||
|
||||
//! the update operator |
||||
void operator()(const oclMat& frame, oclMat& fgmask, float learningRate = 0.f); |
||||
|
||||
//! computes a background image which are the mean of all background gaussians |
||||
void getBackgroundImage(oclMat& backgroundImage) const; |
||||
|
||||
//! releases all inner buffers |
||||
void release(); |
||||
|
||||
int history; |
||||
float varThreshold; |
||||
float backgroundRatio; |
||||
float noiseSigma; |
||||
|
||||
private: |
||||
/* hidden */ |
||||
}; |
||||
|
||||
The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [MOG2001]_. |
||||
|
||||
.. seealso:: :ocv:class:`BackgroundSubtractorMOG` |
||||
|
||||
|
||||
ocl::MOG::MOG |
||||
--------------------- |
||||
The constructor. |
||||
|
||||
.. ocv:function:: ocl::MOG::MOG(int nmixtures = -1) |
||||
|
||||
:param nmixtures: Number of Gaussian mixtures. |
||||
|
||||
Default constructor sets all parameters to default values. |
||||
|
||||
|
||||
ocl::MOG::operator() |
||||
------------------------ |
||||
Updates the background model and returns the foreground mask. |
||||
|
||||
.. ocv:function:: void ocl::MOG::operator()(const oclMat& frame, oclMat& fgmask, float learningRate = 0.f) |
||||
|
||||
:param frame: Next video frame. |
||||
|
||||
:param fgmask: The output foreground mask as an 8-bit binary image. |
||||
|
||||
:param stream: Stream for the asynchronous version. |
||||
|
||||
|
||||
ocl::MOG::getBackgroundImage |
||||
-------------------------------- |
||||
Computes a background image. |
||||
|
||||
.. ocv:function:: void ocl::MOG::getBackgroundImage(oclMat& backgroundImage) const |
||||
|
||||
:param backgroundImage: The output background image. |
||||
|
||||
:param stream: Stream for the asynchronous version. |
||||
|
||||
|
||||
ocl::MOG::release |
||||
--------------------- |
||||
Releases all inner buffer's memory. |
||||
|
||||
.. ocv:function:: void ocl::MOG::release() |
||||
|
||||
|
||||
ocl::MOG2 |
||||
------------- |
||||
.. ocv:class:: ocl::MOG2 : public ocl::BackgroundSubtractor |
||||
|
||||
Gaussian Mixture-based Background/Foreground Segmentation Algorithm. :: |
||||
|
||||
class CV_EXPORTS MOG2: public cv::ocl::BackgroundSubtractor |
||||
{ |
||||
public: |
||||
//! the default constructor |
||||
MOG2(int nmixtures = -1); |
||||
|
||||
//! re-initiaization method |
||||
void initialize(Size frameSize, int frameType); |
||||
|
||||
//! the update operator |
||||
void operator()(const oclMat& frame, oclMat& fgmask, float learningRate = -1.0f); |
||||
|
||||
//! computes a background image which are the mean of all background gaussians |
||||
void getBackgroundImage(oclMat& backgroundImage) const; |
||||
|
||||
//! releases all inner buffers |
||||
void release(); |
||||
|
||||
int history; |
||||
|
||||
float varThreshold; |
||||
|
||||
float backgroundRatio; |
||||
|
||||
float varThresholdGen; |
||||
|
||||
float fVarInit; |
||||
float fVarMin; |
||||
float fVarMax; |
||||
|
||||
float fCT; |
||||
|
||||
bool bShadowDetection; |
||||
unsigned char nShadowDetection; |
||||
float fTau; |
||||
|
||||
private: |
||||
/* hidden */ |
||||
}; |
||||
|
||||
The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [MOG2004]_. |
||||
|
||||
Here are important members of the class that control the algorithm, which you can set after constructing the class instance: |
||||
|
||||
.. ocv:member:: float backgroundRatio |
||||
|
||||
Threshold defining whether the component is significant enough to be included into the background model. ``cf=0.1 => TB=0.9`` is default. For ``alpha=0.001``, it means that the mode should exist for approximately 105 frames before it is considered foreground. |
||||
|
||||
.. ocv:member:: float varThreshold |
||||
|
||||
Threshold for the squared Mahalanobis distance that helps decide when a sample is close to the existing components (corresponds to ``Tg``). If it is not close to any component, a new component is generated. ``3 sigma => Tg=3*3=9`` is default. A smaller ``Tg`` value generates more components. A higher ``Tg`` value may result in a small number of components but they can grow too large. |
||||
|
||||
.. ocv:member:: float fVarInit |
||||
|
||||
Initial variance for the newly generated components. It affects the speed of adaptation. The parameter value is based on your estimate of the typical standard deviation from the images. OpenCV uses 15 as a reasonable value. |
||||
|
||||
.. ocv:member:: float fVarMin |
||||
|
||||
Parameter used to further control the variance. |
||||
|
||||
.. ocv:member:: float fVarMax |
||||
|
||||
Parameter used to further control the variance. |
||||
|
||||
.. ocv:member:: float fCT |
||||
|
||||
Complexity reduction parameter. This parameter defines the number of samples needed to accept to prove the component exists. ``CT=0.05`` is a default value for all the samples. By setting ``CT=0`` you get an algorithm very similar to the standard Stauffer&Grimson algorithm. |
||||
|
||||
.. ocv:member:: uchar nShadowDetection |
||||
|
||||
The value for marking shadow pixels in the output foreground mask. Default value is 127. |
||||
|
||||
.. ocv:member:: float fTau |
||||
|
||||
Shadow threshold. The shadow is detected if the pixel is a darker version of the background. ``Tau`` is a threshold defining how much darker the shadow can be. ``Tau= 0.5`` means that if a pixel is more than twice darker then it is not shadow. See [ShadowDetect2003]_. |
||||
|
||||
.. ocv:member:: bool bShadowDetection |
||||
|
||||
Parameter defining whether shadow detection should be enabled. |
||||
|
||||
.. seealso:: :ocv:class:`BackgroundSubtractorMOG2` |
||||
|
||||
|
||||
ocl::MOG2::MOG2 |
||||
----------------------- |
||||
The constructor. |
||||
|
||||
.. ocv:function:: ocl::MOG2::MOG2(int nmixtures = -1) |
||||
|
||||
:param nmixtures: Number of Gaussian mixtures. |
||||
|
||||
Default constructor sets all parameters to default values. |
||||
|
||||
|
||||
ocl::MOG2::operator() |
||||
------------------------- |
||||
Updates the background model and returns the foreground mask. |
||||
|
||||
.. ocv:function:: void ocl::MOG2::operator()( const oclMat& frame, oclMat& fgmask, float learningRate=-1.0f) |
||||
|
||||
:param frame: Next video frame. |
||||
|
||||
:param fgmask: The output foreground mask as an 8-bit binary image. |
||||
|
||||
:param stream: Stream for the asynchronous version. |
||||
|
||||
|
||||
ocl::MOG2::getBackgroundImage |
||||
--------------------------------- |
||||
Computes a background image. |
||||
|
||||
.. ocv:function:: void ocl::MOG2::getBackgroundImage(oclMat& backgroundImage) const |
||||
|
||||
:param backgroundImage: The output background image. |
||||
|
||||
:param stream: Stream for the asynchronous version. |
||||
|
||||
|
||||
ocl::MOG2::release |
||||
---------------------- |
||||
Releases all inner buffer's memory. |
||||
|
||||
.. ocv:function:: void ocl::MOG2::release() |
||||
@ -0,0 +1,93 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Fangfang Bai, fangfang@multicorewareinc.com
|
||||
// Jin Ma, jin@multicorewareinc.com
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||
// this list of conditions and the following disclaimer in the documentation
|
||||
// and/or other oclMaterials provided with the distribution.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// This software is provided by the copyright holders and contributors as is and
|
||||
// any express or implied warranties, including, but not limited to, the implied
|
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||
// indirect, incidental, special, exemplary, or consequential damages
|
||||
// (including, but not limited to, procurement of substitute goods or services;
|
||||
// loss of use, data, or profits; or business interruption) however caused
|
||||
// and on any theory of liability, whether in contract, strict liability,
|
||||
// or tort (including negligence or otherwise) arising in any way out of
|
||||
// the use of this software, even if advised of the possibility of such damage.
|
||||
//
|
||||
//M*/
|
||||
#include "perf_precomp.hpp" |
||||
using namespace perf; |
||||
using namespace std; |
||||
using namespace cv::ocl; |
||||
using namespace cv; |
||||
using std::tr1::tuple; |
||||
using std::tr1::get; |
||||
///////////// Kalman Filter ////////////////////////
|
||||
|
||||
typedef tuple<int> KalmanFilterType; |
||||
typedef TestBaseWithParam<KalmanFilterType> KalmanFilterFixture; |
||||
|
||||
PERF_TEST_P(KalmanFilterFixture, KalmanFilter, |
||||
::testing::Values(1000, 1500)) |
||||
{ |
||||
KalmanFilterType params = GetParam(); |
||||
const int dim = get<0>(params); |
||||
|
||||
cv::Mat sample(dim, 1, CV_32FC1), dresult; |
||||
randu(sample, -1, 1); |
||||
|
||||
cv::Mat statePre_; |
||||
|
||||
if(RUN_PLAIN_IMPL) |
||||
{ |
||||
cv::KalmanFilter kalman; |
||||
TEST_CYCLE() |
||||
{ |
||||
kalman.init(dim, dim); |
||||
kalman.correct(sample); |
||||
kalman.predict(); |
||||
} |
||||
statePre_ = kalman.statePre; |
||||
}else if(RUN_OCL_IMPL) |
||||
{ |
||||
cv::ocl::oclMat dsample(sample); |
||||
cv::ocl::KalmanFilter kalman_ocl; |
||||
OCL_TEST_CYCLE() |
||||
{ |
||||
kalman_ocl.init(dim, dim); |
||||
kalman_ocl.correct(dsample); |
||||
kalman_ocl.predict(); |
||||
} |
||||
kalman_ocl.statePre.download(statePre_); |
||||
}else |
||||
OCL_PERF_ELSE |
||||
SANITY_CHECK(statePre_); |
||||
} |
||||
@ -0,0 +1,109 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Jin Ma, jin@multicorewareinc.com
|
||||
// Xiaopeng Fu, fuxiaopeng2222@163.com
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||
// this list of conditions and the following disclaimer in the documentation
|
||||
// and/or other oclMaterials provided with the distribution.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// This software is provided by the copyright holders and contributors as is and
|
||||
// any express or implied warranties, including, but not limited to, the implied
|
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||
// indirect, incidental, special, exemplary, or consequential damages
|
||||
// (including, but not limited to, procurement of substitute goods or services;
|
||||
// loss of use, data, or profits; or business interruption) however caused
|
||||
// and on any theory of liability, whether in contract, strict liability,
|
||||
// or tort (including negligence or otherwise) arising in any way out of
|
||||
// the use of this software, even if advised of the possibility of such damage.
|
||||
//
|
||||
//M*/
|
||||
#include "perf_precomp.hpp" |
||||
using namespace perf; |
||||
using namespace std; |
||||
using namespace cv::ocl; |
||||
using namespace cv; |
||||
using std::tr1::tuple; |
||||
using std::tr1::get; |
||||
////////////////////////////////// K-NEAREST NEIGHBOR ////////////////////////////////////
|
||||
static void genData(Mat& trainData, Size size, Mat& trainLabel = Mat().setTo(Scalar::all(0)), int nClasses = 0) |
||||
{ |
||||
trainData.create(size, CV_32FC1); |
||||
randu(trainData, 1.0, 100.0); |
||||
|
||||
if(nClasses != 0) |
||||
{ |
||||
trainLabel.create(size.height, 1, CV_8UC1); |
||||
randu(trainLabel, 0, nClasses - 1); |
||||
trainLabel.convertTo(trainLabel, CV_32FC1); |
||||
} |
||||
} |
||||
|
||||
typedef tuple<int> KNNParamType; |
||||
typedef TestBaseWithParam<KNNParamType> KNNFixture; |
||||
|
||||
PERF_TEST_P(KNNFixture, KNN, |
||||
testing::Values(1000, 2000, 4000)) |
||||
{ |
||||
KNNParamType params = GetParam(); |
||||
const int rows = get<0>(params); |
||||
int columns = 100; |
||||
int k = rows/250; |
||||
|
||||
Mat trainData, trainLabels; |
||||
Size size(columns, rows); |
||||
genData(trainData, size, trainLabels, 3); |
||||
|
||||
Mat testData; |
||||
genData(testData, size); |
||||
Mat best_label; |
||||
|
||||
if(RUN_PLAIN_IMPL) |
||||
{ |
||||
TEST_CYCLE() |
||||
{ |
||||
CvKNearest knn_cpu; |
||||
knn_cpu.train(trainData, trainLabels); |
||||
knn_cpu.find_nearest(testData, k, &best_label); |
||||
} |
||||
}else if(RUN_OCL_IMPL) |
||||
{ |
||||
cv::ocl::oclMat best_label_ocl; |
||||
cv::ocl::oclMat testdata; |
||||
testdata.upload(testData); |
||||
|
||||
OCL_TEST_CYCLE() |
||||
{ |
||||
cv::ocl::KNearestNeighbour knn_ocl; |
||||
knn_ocl.train(trainData, trainLabels); |
||||
knn_ocl.find_nearest(testdata, k, best_label_ocl); |
||||
} |
||||
best_label_ocl.download(best_label); |
||||
}else |
||||
OCL_PERF_ELSE |
||||
SANITY_CHECK(best_label); |
||||
} |
||||
File diff suppressed because it is too large
Load Diff
@ -0,0 +1,135 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Jin Ma, jin@multicorewareinc.com
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||
// this list of conditions and the following disclaimer in the documentation
|
||||
// and/or other oclMaterials provided with the distribution.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// This software is provided by the copyright holders and contributors as is and
|
||||
// any express or implied warranties, including, but not limited to, the implied
|
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||
// indirect, incidental, special, exemplary, or consequential damages
|
||||
// (including, but not limited to, procurement of substitute goods or services;
|
||||
// loss of use, data, or profits; or business interruption) however caused
|
||||
// and on any theory of liability, whether in contract, strict liability,
|
||||
// or tort (including negligence or otherwise) arising in any way out of
|
||||
// the use of this software, even if advised of the possibility of such damage.
|
||||
//
|
||||
//M*/
|
||||
#include "precomp.hpp" |
||||
|
||||
using namespace std; |
||||
using namespace cv; |
||||
using namespace cv::ocl; |
||||
|
||||
KalmanFilter::KalmanFilter() |
||||
{ |
||||
|
||||
} |
||||
|
||||
KalmanFilter::KalmanFilter(int dynamParams, int measureParams, int controlParams, int type) |
||||
{ |
||||
init(dynamParams, measureParams, controlParams, type); |
||||
} |
||||
|
||||
void KalmanFilter::init(int DP, int MP, int CP, int type) |
||||
{ |
||||
CV_Assert( DP > 0 && MP > 0 ); |
||||
CV_Assert( type == CV_32F || type == CV_64F ); |
||||
CP = cv::max(CP, 0); |
||||
|
||||
statePre.create(DP, 1, type); |
||||
statePre.setTo(Scalar::all(0)); |
||||
|
||||
statePost.create(DP, 1, type); |
||||
statePost.setTo(Scalar::all(0)); |
||||
|
||||
transitionMatrix.create(DP, DP, type); |
||||
setIdentity(transitionMatrix, 1); |
||||
|
||||
processNoiseCov.create(DP, DP, type); |
||||
setIdentity(processNoiseCov, 1); |
||||
|
||||
measurementNoiseCov.create(MP, MP, type); |
||||
setIdentity(measurementNoiseCov, 1); |
||||
|
||||
measurementMatrix.create(MP, DP, type); |
||||
measurementMatrix.setTo(Scalar::all(0)); |
||||
|
||||
errorCovPre.create(DP, DP, type); |
||||
errorCovPre.setTo(Scalar::all(0)); |
||||
|
||||
errorCovPost.create(DP, DP, type); |
||||
errorCovPost.setTo(Scalar::all(0)); |
||||
|
||||
gain.create(DP, MP, type); |
||||
gain.setTo(Scalar::all(0)); |
||||
|
||||
if( CP > 0 ) |
||||
{ |
||||
controlMatrix.create(DP, CP, type); |
||||
controlMatrix.setTo(Scalar::all(0)); |
||||
} |
||||
else |
||||
controlMatrix.release(); |
||||
|
||||
temp1.create(DP, DP, type); |
||||
temp2.create(MP, DP, type); |
||||
temp3.create(MP, MP, type); |
||||
temp4.create(MP, DP, type); |
||||
temp5.create(MP, 1, type); |
||||
} |
||||
|
||||
CV_EXPORTS const oclMat& KalmanFilter::predict(const oclMat& control) |
||||
{ |
||||
gemm(transitionMatrix, statePost, 1, oclMat(), 0, statePre); |
||||
oclMat temp; |
||||
|
||||
if(control.data) |
||||
gemm(controlMatrix, control, 1, statePre, 1, statePre); |
||||
gemm(transitionMatrix, errorCovPost, 1, oclMat(), 0, temp1); |
||||
gemm(temp1, transitionMatrix, 1, processNoiseCov, 1, errorCovPre, GEMM_2_T); |
||||
statePre.copyTo(statePost); |
||||
return statePre; |
||||
} |
||||
|
||||
CV_EXPORTS const oclMat& KalmanFilter::correct(const oclMat& measurement) |
||||
{ |
||||
CV_Assert(measurement.empty() == false); |
||||
gemm(measurementMatrix, errorCovPre, 1, oclMat(), 0, temp2); |
||||
gemm(temp2, measurementMatrix, 1, measurementNoiseCov, 1, temp3, GEMM_2_T); |
||||
Mat temp; |
||||
solve(Mat(temp3), Mat(temp2), temp, DECOMP_SVD); |
||||
temp4.upload(temp); |
||||
gain = temp4.t(); |
||||
gemm(measurementMatrix, statePre, -1, measurement, 1, temp5); |
||||
gemm(gain, temp5, 1, statePre, 1, statePost); |
||||
gemm(gain, temp2, -1, errorCovPre, 1, errorCovPost); |
||||
return statePost; |
||||
} |
||||
@ -0,0 +1,157 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Jin Ma, jin@multicorewareinc.com
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||
// this list of conditions and the following disclaimer in the documentation
|
||||
// and/or other oclMaterials provided with the distribution.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// This software is provided by the copyright holders and contributors "as is" and
|
||||
// any express or implied warranties, including, but not limited to, the implied
|
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||
// indirect, incidental, special, exemplary, or consequential damages
|
||||
// (including, but not limited to, procurement of substitute goods or services;
|
||||
// loss of use, data, or profits; or business interruption) however caused
|
||||
// and on any theory of liability, whether in contract, strict liability,
|
||||
// or tort (including negligence or otherwise) arising in any way out of
|
||||
// the use of this software, even if advised of the possibility of such damage.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#include "precomp.hpp" |
||||
using namespace cv; |
||||
using namespace cv::ocl; |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace ocl |
||||
{ |
||||
extern const char* knearest;//knearest
|
||||
} |
||||
} |
||||
|
||||
KNearestNeighbour::KNearestNeighbour() |
||||
{ |
||||
clear(); |
||||
} |
||||
|
||||
KNearestNeighbour::~KNearestNeighbour() |
||||
{ |
||||
clear(); |
||||
samples_ocl.release(); |
||||
} |
||||
|
||||
void KNearestNeighbour::clear() |
||||
{ |
||||
CvKNearest::clear(); |
||||
} |
||||
|
||||
bool KNearestNeighbour::train(const Mat& trainData, Mat& labels, Mat& sampleIdx, |
||||
bool isRegression, int _max_k, bool updateBase) |
||||
{ |
||||
max_k = _max_k; |
||||
bool cv_knn_train = CvKNearest::train(trainData, labels, sampleIdx, isRegression, max_k, updateBase); |
||||
|
||||
CvVectors* s = CvKNearest::samples; |
||||
|
||||
cv::Mat samples_mat(s->count, CvKNearest::var_count + 1, s->type); |
||||
|
||||
float* s1 = (float*)(s + 1); |
||||
for(int i = 0; i < s->count; i++) |
||||
{ |
||||
float* t1 = s->data.fl[i]; |
||||
for(int j = 0; j < CvKNearest::var_count; j++) |
||||
{ |
||||
Point pos(j, i); |
||||
samples_mat.at<float>(pos) = t1[j]; |
||||
} |
||||
|
||||
Point pos_label(CvKNearest::var_count, i); |
||||
samples_mat.at<float>(pos_label) = s1[i]; |
||||
} |
||||
|
||||
samples_ocl = samples_mat; |
||||
return cv_knn_train; |
||||
} |
||||
|
||||
void KNearestNeighbour::find_nearest(const oclMat& samples, int k, oclMat& lables) |
||||
{ |
||||
CV_Assert(!samples_ocl.empty()); |
||||
lables.create(samples.rows, 1, CV_32FC1); |
||||
|
||||
CV_Assert(samples.cols == CvKNearest::var_count); |
||||
CV_Assert(samples.type() == CV_32FC1); |
||||
CV_Assert(k >= 1 && k <= max_k); |
||||
|
||||
int k1 = KNearest::get_sample_count(); |
||||
k1 = MIN( k1, k ); |
||||
|
||||
String kernel_name = "knn_find_nearest"; |
||||
cl_ulong local_memory_size = queryLocalMemInfo(); |
||||
int nThreads = local_memory_size / (2 * k * 4); |
||||
if(nThreads >= 256) |
||||
nThreads = 256; |
||||
|
||||
int smem_size = nThreads * k * 4 * 2; |
||||
size_t local_thread[] = {1, nThreads, 1}; |
||||
size_t global_thread[] = {1, samples.rows, 1}; |
||||
|
||||
char build_option[50]; |
||||
if(!Context::getContext()->supportsFeature(Context::CL_DOUBLE)) |
||||
{ |
||||
sprintf(build_option, " "); |
||||
}else |
||||
sprintf(build_option, "-D DOUBLE_SUPPORT"); |
||||
|
||||
std::vector< std::pair<size_t, const void*> > args; |
||||
|
||||
int samples_ocl_step = samples_ocl.step/samples_ocl.elemSize(); |
||||
int samples_step = samples.step/samples.elemSize(); |
||||
int lables_step = lables.step/lables.elemSize(); |
||||
|
||||
int _regression = 0; |
||||
if(CvKNearest::regression) |
||||
_regression = 1; |
||||
|
||||
args.push_back(std::make_pair(sizeof(cl_mem), (void*)&samples.data)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&samples.rows)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&samples.cols)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&samples_step)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&k)); |
||||
args.push_back(std::make_pair(sizeof(cl_mem), (void*)&samples_ocl.data)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&samples_ocl.rows)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&samples_ocl_step)); |
||||
args.push_back(std::make_pair(sizeof(cl_mem), (void*)&lables.data)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&lables_step)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&_regression)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&k1)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&samples_ocl.cols)); |
||||
args.push_back(std::make_pair(sizeof(cl_int), (void*)&nThreads)); |
||||
args.push_back(std::make_pair(smem_size, (void*)NULL)); |
||||
openCLExecuteKernel(Context::getContext(), &knearest, kernel_name, global_thread, local_thread, args, -1, -1, build_option); |
||||
} |
||||
@ -1,158 +0,0 @@ |
||||
/*M/////////////////////////////////////////////////////////////////////////////////////// |
||||
// |
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. |
||||
// |
||||
// By downloading, copying, installing or using the software you agree to this license. |
||||
// If you do not agree to this license, do not download, install, |
||||
// copy or use the software. |
||||
// |
||||
// |
||||
// License Agreement |
||||
// For Open Source Computer Vision Library |
||||
// |
||||
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. |
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. |
||||
// Third party copyrights are property of their respective owners. |
||||
// |
||||
// @Authors |
||||
// Shengen Yan,yanshengen@gmail.com |
||||
// |
||||
// Redistribution and use in source and binary forms, with or without modification, |
||||
// are permitted provided that the following conditions are met: |
||||
// |
||||
// * Redistribution's of source code must retain the above copyright notice, |
||||
// this list of conditions and the following disclaimer. |
||||
// |
||||
// * Redistribution's in binary form must reproduce the above copyright notice, |
||||
// this list of conditions and the following disclaimer in the documentation |
||||
// and/or other oclMaterials provided with the distribution. |
||||
// |
||||
// * The name of the copyright holders may not be used to endorse or promote products |
||||
// derived from this software without specific prior written permission. |
||||
// |
||||
// This software is provided by the copyright holders and contributors as is and |
||||
// any express or implied warranties, including, but not limited to, the implied |
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed. |
||||
// In no event shall the Intel Corporation or contributors be liable for any direct, |
||||
// indirect, incidental, special, exemplary, or consequential damages |
||||
// (including, but not limited to, procurement of substitute goods or services; |
||||
// loss of use, data, or profits; or business interruption) however caused |
||||
// and on any theory of liability, whether in contract, strict liability, |
||||
// or tort (including negligence or otherwise) arising in any way out of |
||||
// the use of this software, even if advised of the possibility of such damage. |
||||
// |
||||
//M*/ |
||||
|
||||
/**************************************PUBLICFUNC*************************************/ |
||||
#if defined (DOUBLE_SUPPORT) |
||||
#pragma OPENCL EXTENSION cl_khr_fp64:enable |
||||
#endif |
||||
|
||||
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics:enable |
||||
#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics:enable |
||||
#define CV_PI 3.1415926535897932384626433832795 |
||||
|
||||
char round_char(double v){ |
||||
char v1=(char)v; |
||||
return convert_char_sat(v+(v>=0 ? 0.5 : -0.5)); |
||||
} |
||||
unsigned char round_uchar(double v){ |
||||
unsigned char v1=(unsigned char)v; |
||||
return convert_uchar_sat(v+(v>=0 ? 0.5 : -0.5)); |
||||
} |
||||
short round_short(double v){ |
||||
short v1=(short)v; |
||||
return convert_short_sat(v+(v>=0 ? 0.5 : -0.5)); |
||||
} |
||||
unsigned short round_ushort(double v){ |
||||
unsigned short v1=(unsigned short)v; |
||||
return convert_ushort_sat(v+(v>=0 ? 0.5 : -0.5)); |
||||
} |
||||
int round_int(double v){ |
||||
int v1=(int)v; |
||||
return convert_int_sat(v+(v>=0 ? 0.5 : -0.5)); |
||||
} |
||||
|
||||
char round2_char(double v){ |
||||
char v1=(char)v; |
||||
if((v-v1)==0.5&&v1%2==0) |
||||
return v1; |
||||
else |
||||
return convert_char_sat(v+(v>=0 ? 0.5 : -0.5)); |
||||
} |
||||
unsigned char round2_uchar(double v){ |
||||
unsigned char v1=(unsigned char)v; |
||||
if((v-v1)==0.5&&v1%2==0) |
||||
return v1; |
||||
else |
||||
return convert_uchar_sat(v+(v>=0 ? 0.5 : -0.5)); |
||||
} |
||||
short round2_short(double v){ |
||||
short v1=(short)v; |
||||
if((v-v1)==0.5&&v1%2==0) |
||||
return v1; |
||||
else |
||||
return convert_short_sat(v+(v>=0 ? 0.5 : -0.5)); |
||||
} |
||||
unsigned short round2_ushort(double v){ |
||||
unsigned short v1=(unsigned short)v; |
||||
if((v-v1)==0.5&&v1%2==0) |
||||
return v1; |
||||
else |
||||
return convert_ushort_sat(v+(v>=0 ? 0.5 : -0.5)); |
||||
} |
||||
int round2_int(double v){ |
||||
int v1=(int)v; |
||||
if((v-v1)==0.5&&v1%2==0) |
||||
return v1; |
||||
else |
||||
return convert_int_sat(v+(v>=0 ? 0.5 : -0.5)); |
||||
} |
||||
|
||||
/*****************************************EXP***************************************/ |
||||
__kernel void arithm_op_exp_5 (int rows,int cols,int srcStep,__global float *src1Mat, |
||||
__global float * dstMat,int channels) |
||||
{ |
||||
size_t x = get_global_id(0); |
||||
size_t y = get_global_id(1); |
||||
if (x < cols && y < rows) |
||||
{ |
||||
size_t idx = y * ( srcStep >> 2 ) + x; |
||||
dstMat[idx] = (float)exp((float)src1Mat[idx]); |
||||
} |
||||
} |
||||
__kernel void arithm_op_exp_6 (int rows,int cols,int srcStep,__global double *src1Mat, |
||||
__global double * dstMat,int channels) |
||||
{ |
||||
size_t x = get_global_id(0); |
||||
size_t y = get_global_id(1); |
||||
if (x < cols && y < rows) |
||||
{ |
||||
size_t idx = y * ( srcStep >> 3 ) + x; |
||||
dstMat[idx] = exp(src1Mat[idx]); |
||||
} |
||||
} |
||||
|
||||
/*****************************************LOG***************************************/ |
||||
__kernel void arithm_op_log_5 (int rows,int cols,int srcStep,__global float *src1Mat, |
||||
__global float * dstMat,int channels) |
||||
{ |
||||
size_t x = get_global_id(0); |
||||
size_t y = get_global_id(1); |
||||
if (x < cols && y < rows) |
||||
{ |
||||
size_t idx = y * ( srcStep >> 2 ) + x; |
||||
dstMat[idx] =(float) log((float)src1Mat[idx]); |
||||
} |
||||
} |
||||
__kernel void arithm_op_log_6 (int rows,int cols,int srcStep,__global double *src1Mat, |
||||
__global double * dstMat,int channels) |
||||
{ |
||||
size_t x = get_global_id(0); |
||||
size_t y = get_global_id(1); |
||||
if (x < cols && y < rows) |
||||
{ |
||||
size_t idx = y * ( srcStep >> 3 ) + x; |
||||
dstMat[idx] = log(src1Mat[idx]); |
||||
} |
||||
} |
||||
@ -1,970 +0,0 @@ |
||||
/*M/////////////////////////////////////////////////////////////////////////////////////// |
||||
// |
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. |
||||
// |
||||
// By downloading, copying, installing or using the software you agree to this license. |
||||
// If you do not agree to this license, do not download, install, |
||||
// copy or use the software. |
||||
// |
||||
// |
||||
// License Agreement |
||||
// For Open Source Computer Vision Library |
||||
// |
||||
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. |
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. |
||||
// Third party copyrights are property of their respective owners. |
||||
// |
||||
// @Authors |
||||
// Jia Haipeng, jiahaipeng95@gmail.com |
||||
// |
||||
// Redistribution and use in source and binary forms, with or without modification, |
||||
// are permitted provided that the following conditions are met: |
||||
// |
||||
// * Redistribution's of source code must retain the above copyright notice, |
||||
// this list of conditions and the following disclaimer. |
||||
// |
||||
// * Redistribution's in binary form must reproduce the above copyright notice, |
||||
// this list of conditions and the following disclaimer in the documentation |
||||
// and/or other oclMaterials provided with the distribution. |
||||
// |
||||
// * The name of the copyright holders may not be used to endorse or promote products |
||||
// derived from this software without specific prior written permission. |
||||
// |
||||
// This software is provided by the copyright holders and contributors as is and |
||||
// any express or implied warranties, including, but not limited to, the implied |
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed. |
||||
// In no event shall the Intel Corporation or contributors be liable for any direct, |
||||
// indirect, incidental, special, exemplary, or consequential damages |
||||
// (including, but not limited to, procurement of substitute goods or services; |
||||
// loss of use, data, or profits; or business interruption) however caused |
||||
// and on any theory of liability, whether in contract, strict liability, |
||||
// or tort (including negligence or otherwise) arising in any way out of |
||||
// the use of this software, even if advised of the possibility of such damage. |
||||
// |
||||
//M*/ |
||||
|
||||
#if defined (DOUBLE_SUPPORT) |
||||
#ifdef cl_khr_fp64 |
||||
#pragma OPENCL EXTENSION cl_khr_fp64:enable |
||||
#elif defined (cl_amd_fp64) |
||||
#pragma OPENCL EXTENSION cl_amd_fp64:enable |
||||
#endif |
||||
#endif |
||||
|
||||
////////////////////////////////////////////////////////////////////////////////////////////////////// |
||||
/////////////////////////////////////////////absdiff//////////////////////////////////////////////////// |
||||
/////////////////////////////////////////////////////////////////////////////////////////////////////// |
||||
/**************************************adddiff *************************************/ |
||||
__kernel void arithm_absdiff_D0 (__global uchar *src1, int src1_step, int src1_offset, |
||||
__global uchar *src2, int src2_step, int src2_offset, |
||||
__global uchar *dst, int dst_step, int dst_offset, |
||||
int rows, int cols, int dst_step1) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
x = x << 2; |
||||
|
||||
#ifdef dst_align |
||||
#undef dst_align |
||||
#endif |
||||
#define dst_align (dst_offset & 3) |
||||
int src1_index = mad24(y, src1_step, x + src1_offset - dst_align); |
||||
int src2_index = mad24(y, src2_step, x + src2_offset - dst_align); |
||||
|
||||
int dst_start = mad24(y, dst_step, dst_offset); |
||||
int dst_end = mad24(y, dst_step, dst_offset + dst_step1); |
||||
int dst_index = mad24(y, dst_step, dst_offset + x & (int)0xfffffffc); |
||||
int src1_index_fix = src1_index < 0 ? 0 : src1_index; |
||||
int src2_index_fix = src2_index < 0 ? 0 : src2_index; |
||||
uchar4 src1_data = vload4(0, src1 + src1_index_fix); |
||||
uchar4 src2_data = vload4(0, src2 + src2_index_fix); |
||||
if(src1_index < 0) |
||||
{ |
||||
uchar4 tmp; |
||||
tmp.xyzw = (src1_index == -2) ? src1_data.zwxy:src1_data.yzwx; |
||||
src1_data.xyzw = (src1_index == -1) ? src1_data.wxyz:tmp.xyzw; |
||||
} |
||||
if(src2_index < 0) |
||||
{ |
||||
uchar4 tmp; |
||||
tmp.xyzw = (src2_index == -2) ? src2_data.zwxy:src2_data.yzwx; |
||||
src2_data.xyzw = (src2_index == -1) ? src2_data.wxyz:tmp.xyzw; |
||||
} |
||||
|
||||
uchar4 dst_data = *((__global uchar4 *)(dst + dst_index)); |
||||
uchar4 tmp_data = abs_diff(src1_data, src2_data); |
||||
|
||||
dst_data.x = ((dst_index + 0 >= dst_start) && (dst_index + 0 < dst_end)) ? tmp_data.x : dst_data.x; |
||||
dst_data.y = ((dst_index + 1 >= dst_start) && (dst_index + 1 < dst_end)) ? tmp_data.y : dst_data.y; |
||||
dst_data.z = ((dst_index + 2 >= dst_start) && (dst_index + 2 < dst_end)) ? tmp_data.z : dst_data.z; |
||||
dst_data.w = ((dst_index + 3 >= dst_start) && (dst_index + 3 < dst_end)) ? tmp_data.w : dst_data.w; |
||||
|
||||
*((__global uchar4 *)(dst + dst_index)) = dst_data; |
||||
} |
||||
} |
||||
__kernel void arithm_absdiff_D2 (__global ushort *src1, int src1_step, int src1_offset, |
||||
__global ushort *src2, int src2_step, int src2_offset, |
||||
__global ushort *dst, int dst_step, int dst_offset, |
||||
int rows, int cols, int dst_step1) |
||||
|
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
x = x << 2; |
||||
|
||||
#ifdef dst_align |
||||
#undef dst_align |
||||
#endif |
||||
#define dst_align ((dst_offset >> 1) & 3) |
||||
int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); |
||||
int src2_index = mad24(y, src2_step, (x << 1) + src2_offset - (dst_align << 1)); |
||||
|
||||
int dst_start = mad24(y, dst_step, dst_offset); |
||||
int dst_end = mad24(y, dst_step, dst_offset + dst_step1); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x << 1) & (int)0xfffffff8); |
||||
|
||||
ushort4 src1_data = vload4(0, (__global ushort *)((__global char *)src1 + src1_index)); |
||||
ushort4 src2_data = vload4(0, (__global ushort *)((__global char *)src2 + src2_index)); |
||||
|
||||
ushort4 dst_data = *((__global ushort4 *)((__global char *)dst + dst_index)); |
||||
ushort4 tmp_data = abs_diff(src1_data, src2_data); |
||||
|
||||
dst_data.x = ((dst_index + 0 >= dst_start) && (dst_index + 0 < dst_end)) ? tmp_data.x : dst_data.x; |
||||
dst_data.y = ((dst_index + 2 >= dst_start) && (dst_index + 2 < dst_end)) ? tmp_data.y : dst_data.y; |
||||
dst_data.z = ((dst_index + 4 >= dst_start) && (dst_index + 4 < dst_end)) ? tmp_data.z : dst_data.z; |
||||
dst_data.w = ((dst_index + 6 >= dst_start) && (dst_index + 6 < dst_end)) ? tmp_data.w : dst_data.w; |
||||
|
||||
*((__global ushort4 *)((__global char *)dst + dst_index)) = dst_data; |
||||
} |
||||
} |
||||
__kernel void arithm_absdiff_D3 (__global short *src1, int src1_step, int src1_offset, |
||||
__global short *src2, int src2_step, int src2_offset, |
||||
__global short *dst, int dst_step, int dst_offset, |
||||
int rows, int cols, int dst_step1) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
x = x << 2; |
||||
|
||||
#ifdef dst_align |
||||
#undef dst_align |
||||
#endif |
||||
#define dst_align ((dst_offset >> 1) & 3) |
||||
int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); |
||||
int src2_index = mad24(y, src2_step, (x << 1) + src2_offset - (dst_align << 1)); |
||||
|
||||
int dst_start = mad24(y, dst_step, dst_offset); |
||||
int dst_end = mad24(y, dst_step, dst_offset + dst_step1); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x << 1) & (int)0xfffffff8); |
||||
|
||||
short4 src1_data = vload4(0, (__global short *)((__global char *)src1 + src1_index)); |
||||
short4 src2_data = vload4(0, (__global short *)((__global char *)src2 + src2_index)); |
||||
|
||||
short4 dst_data = *((__global short4 *)((__global char *)dst + dst_index)); |
||||
ushort4 tmp = abs_diff(src1_data, src2_data); |
||||
short4 tmp_data = convert_short4_sat(tmp); |
||||
|
||||
dst_data.x = ((dst_index + 0 >= dst_start) && (dst_index + 0 < dst_end)) ? tmp_data.x : dst_data.x; |
||||
dst_data.y = ((dst_index + 2 >= dst_start) && (dst_index + 2 < dst_end)) ? tmp_data.y : dst_data.y; |
||||
dst_data.z = ((dst_index + 4 >= dst_start) && (dst_index + 4 < dst_end)) ? tmp_data.z : dst_data.z; |
||||
dst_data.w = ((dst_index + 6 >= dst_start) && (dst_index + 6 < dst_end)) ? tmp_data.w : dst_data.w; |
||||
|
||||
*((__global short4 *)((__global char *)dst + dst_index)) = dst_data; |
||||
} |
||||
} |
||||
|
||||
__kernel void arithm_absdiff_D4 (__global int *src1, int src1_step, int src1_offset, |
||||
__global int *src2, int src2_step, int src2_offset, |
||||
__global int *dst, int dst_step, int dst_offset, |
||||
int rows, int cols, int dst_step1) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 2) + src1_offset); |
||||
int src2_index = mad24(y, src2_step, (x << 2) + src2_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 2) + dst_offset); |
||||
|
||||
int data1 = *((__global int *)((__global char *)src1 + src1_index)); |
||||
int data2 = *((__global int *)((__global char *)src2 + src2_index)); |
||||
uint tmp = abs_diff(data1, data2); |
||||
int tmp_data = convert_int_sat(tmp); |
||||
|
||||
*((__global int *)((__global char *)dst + dst_index)) = tmp_data; |
||||
} |
||||
} |
||||
__kernel void arithm_absdiff_D5 (__global float *src1, int src1_step, int src1_offset, |
||||
__global float *src2, int src2_step, int src2_offset, |
||||
__global float *dst, int dst_step, int dst_offset, |
||||
int rows, int cols, int dst_step1) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 2) + src1_offset); |
||||
int src2_index = mad24(y, src2_step, (x << 2) + src2_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 2) + dst_offset); |
||||
|
||||
float data1 = *((__global float *)((__global char *)src1 + src1_index)); |
||||
float data2 = *((__global float *)((__global char *)src2 + src2_index)); |
||||
float tmp = fabs(data1 - data2); |
||||
|
||||
*((__global float *)((__global char *)dst + dst_index)) = tmp; |
||||
} |
||||
} |
||||
|
||||
#if defined (DOUBLE_SUPPORT) |
||||
__kernel void arithm_absdiff_D6 (__global double *src1, int src1_step, int src1_offset, |
||||
__global double *src2, int src2_step, int src2_offset, |
||||
__global double *dst, int dst_step, int dst_offset, |
||||
int rows, int cols, int dst_step1) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 3) + src1_offset); |
||||
int src2_index = mad24(y, src2_step, (x << 3) + src2_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 3) + dst_offset); |
||||
|
||||
double data1 = *((__global double *)((__global char *)src1 + src1_index)); |
||||
double data2 = *((__global double *)((__global char *)src2 + src2_index)); |
||||
double tmp = fabs(data1-data2); |
||||
|
||||
*((__global double *)((__global char *)dst + dst_index)) = tmp; |
||||
} |
||||
} |
||||
#endif |
||||
|
||||
/**************************************absdiff with scalar**************************************/ |
||||
__kernel void arithm_s_absdiff_C1_D0 (__global uchar *src1, int src1_step, int src1_offset, |
||||
__global uchar *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
x = x << 2; |
||||
|
||||
#ifdef dst_align |
||||
#undef dst_align |
||||
#endif |
||||
#define dst_align (dst_offset & 3) |
||||
int src1_index = mad24(y, src1_step, x + src1_offset - dst_align); |
||||
|
||||
int dst_start = mad24(y, dst_step, dst_offset); |
||||
int dst_end = mad24(y, dst_step, dst_offset + dst_step1); |
||||
int dst_index = mad24(y, dst_step, dst_offset + x & (int)0xfffffffc); |
||||
int src1_index_fix = src1_index < 0 ? 0 : src1_index; |
||||
uchar4 src1_data = vload4(0, src1 + src1_index_fix); |
||||
int4 src2_data = (int4)(src2.x, src2.x, src2.x, src2.x); |
||||
if(src1_index < 0) |
||||
{ |
||||
uchar4 tmp; |
||||
tmp.xyzw = (src1_index == -2) ? src1_data.zwxy:src1_data.yzwx; |
||||
src1_data.xyzw = (src1_index == -1) ? src1_data.wxyz:tmp.xyzw; |
||||
} |
||||
|
||||
uchar4 data = *((__global uchar4 *)(dst + dst_index)); |
||||
uchar4 tmp_data = convert_uchar4_sat(abs_diff(convert_int4_sat(src1_data), src2_data)); |
||||
|
||||
data.x = ((dst_index + 0 >= dst_start) && (dst_index + 0 < dst_end)) ? tmp_data.x : data.x; |
||||
data.y = ((dst_index + 1 >= dst_start) && (dst_index + 1 < dst_end)) ? tmp_data.y : data.y; |
||||
data.z = ((dst_index + 2 >= dst_start) && (dst_index + 2 < dst_end)) ? tmp_data.z : data.z; |
||||
data.w = ((dst_index + 3 >= dst_start) && (dst_index + 3 < dst_end)) ? tmp_data.w : data.w; |
||||
|
||||
*((__global uchar4 *)(dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C1_D2 (__global ushort *src1, int src1_step, int src1_offset, |
||||
__global ushort *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
x = x << 1; |
||||
|
||||
#ifdef dst_align |
||||
#undef dst_align |
||||
#endif |
||||
#define dst_align ((dst_offset >> 1) & 1) |
||||
int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); |
||||
|
||||
int dst_start = mad24(y, dst_step, dst_offset); |
||||
int dst_end = mad24(y, dst_step, dst_offset + dst_step1); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x << 1) & (int)0xfffffffc); |
||||
|
||||
ushort2 src1_data = vload2(0, (__global ushort *)((__global char *)src1 + src1_index)); |
||||
int2 src2_data = (int2)(src2.x, src2.x); |
||||
|
||||
ushort2 data = *((__global ushort2 *)((__global uchar *)dst + dst_index)); |
||||
ushort2 tmp_data = convert_ushort2_sat(abs_diff(convert_int2_sat(src1_data), src2_data)); |
||||
|
||||
data.x = (dst_index + 0 >= dst_start) ? tmp_data.x : data.x; |
||||
data.y = (dst_index + 2 < dst_end ) ? tmp_data.y : data.y; |
||||
|
||||
*((__global ushort2 *)((__global uchar *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C1_D3 (__global short *src1, int src1_step, int src1_offset, |
||||
__global short *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
x = x << 1; |
||||
|
||||
#ifdef dst_align |
||||
#undef dst_align |
||||
#endif |
||||
#define dst_align ((dst_offset >> 1) & 1) |
||||
int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); |
||||
|
||||
int dst_start = mad24(y, dst_step, dst_offset); |
||||
int dst_end = mad24(y, dst_step, dst_offset + dst_step1); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x << 1) & (int)0xfffffffc); |
||||
|
||||
short2 src1_data = vload2(0, (__global short *)((__global char *)src1 + src1_index)); |
||||
int2 src2_data = (int2)(src2.x, src2.x); |
||||
short2 data = *((__global short2 *)((__global uchar *)dst + dst_index)); |
||||
|
||||
ushort2 tmp = convert_ushort2_sat(abs_diff(convert_int2_sat(src1_data), src2_data)); |
||||
short2 tmp_data = convert_short2_sat(tmp); |
||||
|
||||
data.x = (dst_index + 0 >= dst_start) ? tmp_data.x : data.x; |
||||
data.y = (dst_index + 2 < dst_end ) ? tmp_data.y : data.y; |
||||
|
||||
*((__global short2 *)((__global uchar *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C1_D4 (__global int *src1, int src1_step, int src1_offset, |
||||
__global int *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 2) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 2) + dst_offset); |
||||
|
||||
int src_data1 = *((__global int *)((__global char *)src1 + src1_index)); |
||||
int src_data2 = src2.x; |
||||
int dst_data = *((__global int *)((__global char *)dst + dst_index)); |
||||
|
||||
uint tmp_data = abs_diff(src_data1, src_data2); |
||||
int data = convert_int_sat(tmp_data); |
||||
|
||||
*((__global int *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C1_D5 (__global float *src1, int src1_step, int src1_offset, |
||||
__global float *dst, int dst_step, int dst_offset, |
||||
float4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 2) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 2) + dst_offset); |
||||
|
||||
float src_data1 = *((__global float *)((__global char *)src1 + src1_index)); |
||||
float src_data2 = src2.x; |
||||
float dst_data = *((__global float *)((__global char *)dst + dst_index)); |
||||
|
||||
float data = fabs(src_data1 - src_data2); |
||||
|
||||
*((__global float *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
|
||||
#if defined (DOUBLE_SUPPORT) |
||||
__kernel void arithm_s_absdiff_C1_D6 (__global double *src1, int src1_step, int src1_offset, |
||||
__global double *dst, int dst_step, int dst_offset, |
||||
double4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 3) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 3) + dst_offset); |
||||
|
||||
double src_data1 = *((__global double *)((__global char *)src1 + src1_index)); |
||||
double src2_data = src2.x; |
||||
double dst_data = *((__global double *)((__global char *)dst + dst_index)); |
||||
|
||||
double data = fabs(src_data1 - src2_data); |
||||
|
||||
*((__global double *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
#endif |
||||
|
||||
__kernel void arithm_s_absdiff_C2_D0 (__global uchar *src1, int src1_step, int src1_offset, |
||||
__global uchar *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
x = x << 1; |
||||
|
||||
#ifdef dst_align |
||||
#undef dst_align |
||||
#endif |
||||
#define dst_align ((dst_offset >> 1) & 1) |
||||
int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1)); |
||||
|
||||
int dst_start = mad24(y, dst_step, dst_offset); |
||||
int dst_end = mad24(y, dst_step, dst_offset + dst_step1); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x << 1) & (int)0xfffffffc); |
||||
|
||||
uchar4 src1_data = vload4(0, src1 + src1_index); |
||||
int4 src2_data = (int4)(src2.x, src2.y, src2.x, src2.y); |
||||
|
||||
uchar4 data = *((__global uchar4 *)(dst + dst_index)); |
||||
uchar4 tmp_data = convert_uchar4_sat(abs_diff(convert_int4_sat(src1_data), src2_data)); |
||||
|
||||
data.xy = (dst_index + 0 >= dst_start) ? tmp_data.xy : data.xy; |
||||
data.zw = (dst_index + 2 < dst_end ) ? tmp_data.zw : data.zw; |
||||
|
||||
*((__global uchar4 *)(dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C2_D2 (__global ushort *src1, int src1_step, int src1_offset, |
||||
__global ushort *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 2) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 2) + dst_offset); |
||||
|
||||
ushort2 src_data1 = *((__global ushort2 *)((__global char *)src1 + src1_index)); |
||||
int2 src_data2 = (int2)(src2.x, src2.y); |
||||
ushort2 dst_data = *((__global ushort2 *)((__global char *)dst + dst_index)); |
||||
|
||||
ushort2 data = convert_ushort2_sat( abs_diff(convert_int2_sat(src_data1), src_data2)); |
||||
|
||||
*((__global ushort2 *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C2_D3 (__global short *src1, int src1_step, int src1_offset, |
||||
__global short *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 2) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 2) + dst_offset); |
||||
|
||||
short2 src_data1 = *((__global short2 *)((__global char *)src1 + src1_index)); |
||||
int2 src_data2 = (int2)(src2.x, src2.y); |
||||
short2 dst_data = *((__global short2 *)((__global char *)dst + dst_index)); |
||||
|
||||
ushort2 tmp = convert_ushort2_sat(abs_diff(convert_int2_sat(src_data1), src_data2)); |
||||
short2 data = convert_short2_sat(tmp); |
||||
|
||||
*((__global short2 *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C2_D4 (__global int *src1, int src1_step, int src1_offset, |
||||
__global int *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 3) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 3) + dst_offset); |
||||
|
||||
int2 src_data1 = *((__global int2 *)((__global char *)src1 + src1_index)); |
||||
int2 src_data2 = (int2)(src2.x, src2.y); |
||||
int2 dst_data = *((__global int2 *)((__global char *)dst + dst_index)); |
||||
|
||||
int2 data = convert_int2_sat(abs_diff(src_data1, src_data2)); |
||||
*((__global int2 *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C2_D5 (__global float *src1, int src1_step, int src1_offset, |
||||
__global float *dst, int dst_step, int dst_offset, |
||||
float4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 3) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 3) + dst_offset); |
||||
|
||||
float2 src_data1 = *((__global float2 *)((__global char *)src1 + src1_index)); |
||||
float2 src_data2 = (float2)(src2.x, src2.y); |
||||
float2 dst_data = *((__global float2 *)((__global char *)dst + dst_index)); |
||||
|
||||
float2 data = fabs(src_data1 - src_data2); |
||||
*((__global float2 *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
#if defined (DOUBLE_SUPPORT) |
||||
__kernel void arithm_s_absdiff_C2_D6 (__global double *src1, int src1_step, int src1_offset, |
||||
__global double *dst, int dst_step, int dst_offset, |
||||
double4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 4) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 4) + dst_offset); |
||||
|
||||
double2 src_data1 = *((__global double2 *)((__global char *)src1 + src1_index)); |
||||
double2 src_data2 = (double2)(src2.x, src2.y); |
||||
double2 dst_data = *((__global double2 *)((__global char *)dst + dst_index)); |
||||
|
||||
double2 data = fabs(src_data1 - src_data2); |
||||
|
||||
*((__global double2 *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
#endif |
||||
__kernel void arithm_s_absdiff_C3_D0 (__global uchar *src1, int src1_step, int src1_offset, |
||||
__global uchar *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
x = x << 2; |
||||
|
||||
#ifdef dst_align |
||||
#undef dst_align |
||||
#endif |
||||
#define dst_align (((dst_offset % dst_step) / 3 ) & 3) |
||||
int src1_index = mad24(y, src1_step, (x * 3) + src1_offset - (dst_align * 3)); |
||||
|
||||
int dst_start = mad24(y, dst_step, dst_offset); |
||||
int dst_end = mad24(y, dst_step, dst_offset + dst_step1); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x * 3) - (dst_align * 3)); |
||||
|
||||
uchar4 src1_data_0 = vload4(0, src1 + src1_index + 0); |
||||
uchar4 src1_data_1 = vload4(0, src1 + src1_index + 4); |
||||
uchar4 src1_data_2 = vload4(0, src1 + src1_index + 8); |
||||
|
||||
int4 src2_data_0 = (int4)(src2.x, src2.y, src2.z, src2.x); |
||||
int4 src2_data_1 = (int4)(src2.y, src2.z, src2.x, src2.y); |
||||
int4 src2_data_2 = (int4)(src2.z, src2.x, src2.y, src2.z); |
||||
|
||||
uchar4 data_0 = *((__global uchar4 *)(dst + dst_index + 0)); |
||||
uchar4 data_1 = *((__global uchar4 *)(dst + dst_index + 4)); |
||||
uchar4 data_2 = *((__global uchar4 *)(dst + dst_index + 8)); |
||||
|
||||
uchar4 tmp_data_0 = convert_uchar4_sat(abs_diff(convert_int4_sat(src1_data_0), src2_data_0)); |
||||
uchar4 tmp_data_1 = convert_uchar4_sat(abs_diff(convert_int4_sat(src1_data_1), src2_data_1)); |
||||
uchar4 tmp_data_2 = convert_uchar4_sat(abs_diff(convert_int4_sat(src1_data_2), src2_data_2)); |
||||
|
||||
data_0.xyz = ((dst_index + 0 >= dst_start)) ? tmp_data_0.xyz : data_0.xyz; |
||||
data_0.w = ((dst_index + 3 >= dst_start) && (dst_index + 3 < dst_end)) |
||||
? tmp_data_0.w : data_0.w; |
||||
|
||||
data_1.xy = ((dst_index + 3 >= dst_start) && (dst_index + 3 < dst_end)) |
||||
? tmp_data_1.xy : data_1.xy; |
||||
data_1.zw = ((dst_index + 6 >= dst_start) && (dst_index + 6 < dst_end)) |
||||
? tmp_data_1.zw : data_1.zw; |
||||
|
||||
data_2.x = ((dst_index + 6 >= dst_start) && (dst_index + 6 < dst_end)) |
||||
? tmp_data_2.x : data_2.x; |
||||
data_2.yzw = ((dst_index + 9 >= dst_start) && (dst_index + 9 < dst_end)) |
||||
? tmp_data_2.yzw : data_2.yzw; |
||||
|
||||
*((__global uchar4 *)(dst + dst_index + 0)) = data_0; |
||||
*((__global uchar4 *)(dst + dst_index + 4)) = data_1; |
||||
*((__global uchar4 *)(dst + dst_index + 8)) = data_2; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C3_D2 (__global ushort *src1, int src1_step, int src1_offset, |
||||
__global ushort *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
x = x << 1; |
||||
|
||||
#ifdef dst_align |
||||
#undef dst_align |
||||
#endif |
||||
#define dst_align (((dst_offset % dst_step) / 6 ) & 1) |
||||
int src1_index = mad24(y, src1_step, (x * 6) + src1_offset - (dst_align * 6)); |
||||
|
||||
int dst_start = mad24(y, dst_step, dst_offset); |
||||
int dst_end = mad24(y, dst_step, dst_offset + dst_step1); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x * 6) - (dst_align * 6)); |
||||
|
||||
ushort2 src1_data_0 = vload2(0, (__global ushort *)((__global char *)src1 + src1_index + 0)); |
||||
ushort2 src1_data_1 = vload2(0, (__global ushort *)((__global char *)src1 + src1_index + 4)); |
||||
ushort2 src1_data_2 = vload2(0, (__global ushort *)((__global char *)src1 + src1_index + 8)); |
||||
|
||||
int2 src2_data_0 = (int2)(src2.x, src2.y); |
||||
int2 src2_data_1 = (int2)(src2.z, src2.x); |
||||
int2 src2_data_2 = (int2)(src2.y, src2.z); |
||||
|
||||
ushort2 data_0 = *((__global ushort2 *)((__global char *)dst + dst_index + 0)); |
||||
ushort2 data_1 = *((__global ushort2 *)((__global char *)dst + dst_index + 4)); |
||||
ushort2 data_2 = *((__global ushort2 *)((__global char *)dst + dst_index + 8)); |
||||
|
||||
ushort2 tmp_data_0 = convert_ushort2_sat(abs_diff(convert_int2_sat(src1_data_0), src2_data_0)); |
||||
ushort2 tmp_data_1 = convert_ushort2_sat(abs_diff(convert_int2_sat(src1_data_1), src2_data_1)); |
||||
ushort2 tmp_data_2 = convert_ushort2_sat(abs_diff(convert_int2_sat(src1_data_2), src2_data_2)); |
||||
|
||||
data_0.xy = ((dst_index + 0 >= dst_start)) ? tmp_data_0.xy : data_0.xy; |
||||
|
||||
data_1.x = ((dst_index + 0 >= dst_start) && (dst_index + 0 < dst_end)) |
||||
? tmp_data_1.x : data_1.x; |
||||
data_1.y = ((dst_index + 6 >= dst_start) && (dst_index + 6 < dst_end)) |
||||
? tmp_data_1.y : data_1.y; |
||||
|
||||
data_2.xy = ((dst_index + 6 >= dst_start) && (dst_index + 6 < dst_end)) |
||||
? tmp_data_2.xy : data_2.xy; |
||||
|
||||
*((__global ushort2 *)((__global char *)dst + dst_index + 0))= data_0; |
||||
*((__global ushort2 *)((__global char *)dst + dst_index + 4))= data_1; |
||||
*((__global ushort2 *)((__global char *)dst + dst_index + 8))= data_2; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C3_D3 (__global short *src1, int src1_step, int src1_offset, |
||||
__global short *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
x = x << 1; |
||||
|
||||
#ifdef dst_align |
||||
#undef dst_align |
||||
#endif |
||||
#define dst_align (((dst_offset % dst_step) / 6 ) & 1) |
||||
int src1_index = mad24(y, src1_step, (x * 6) + src1_offset - (dst_align * 6)); |
||||
|
||||
int dst_start = mad24(y, dst_step, dst_offset); |
||||
int dst_end = mad24(y, dst_step, dst_offset + dst_step1); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x * 6) - (dst_align * 6)); |
||||
|
||||
short2 src1_data_0 = vload2(0, (__global short *)((__global char *)src1 + src1_index + 0)); |
||||
short2 src1_data_1 = vload2(0, (__global short *)((__global char *)src1 + src1_index + 4)); |
||||
short2 src1_data_2 = vload2(0, (__global short *)((__global char *)src1 + src1_index + 8)); |
||||
|
||||
int2 src2_data_0 = (int2)(src2.x, src2.y); |
||||
int2 src2_data_1 = (int2)(src2.z, src2.x); |
||||
int2 src2_data_2 = (int2)(src2.y, src2.z); |
||||
|
||||
short2 data_0 = *((__global short2 *)((__global char *)dst + dst_index + 0)); |
||||
short2 data_1 = *((__global short2 *)((__global char *)dst + dst_index + 4)); |
||||
short2 data_2 = *((__global short2 *)((__global char *)dst + dst_index + 8)); |
||||
|
||||
short2 tmp_data_0 = convert_short2_sat(abs_diff(convert_int2_sat(src1_data_0), src2_data_0)); |
||||
short2 tmp_data_1 = convert_short2_sat(abs_diff(convert_int2_sat(src1_data_1), src2_data_1)); |
||||
short2 tmp_data_2 = convert_short2_sat(abs_diff(convert_int2_sat(src1_data_2), src2_data_2)); |
||||
|
||||
data_0.xy = ((dst_index + 0 >= dst_start)) ? tmp_data_0.xy : data_0.xy; |
||||
|
||||
data_1.x = ((dst_index + 0 >= dst_start) && (dst_index + 0 < dst_end)) |
||||
? tmp_data_1.x : data_1.x; |
||||
data_1.y = ((dst_index + 6 >= dst_start) && (dst_index + 6 < dst_end)) |
||||
? tmp_data_1.y : data_1.y; |
||||
|
||||
data_2.xy = ((dst_index + 6 >= dst_start) && (dst_index + 6 < dst_end)) |
||||
? tmp_data_2.xy : data_2.xy; |
||||
|
||||
*((__global short2 *)((__global char *)dst + dst_index + 0))= data_0; |
||||
*((__global short2 *)((__global char *)dst + dst_index + 4))= data_1; |
||||
*((__global short2 *)((__global char *)dst + dst_index + 8))= data_2; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C3_D4 (__global int *src1, int src1_step, int src1_offset, |
||||
__global int *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x * 12) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x * 12)); |
||||
|
||||
int src1_data_0 = *((__global int *)((__global char *)src1 + src1_index + 0)); |
||||
int src1_data_1 = *((__global int *)((__global char *)src1 + src1_index + 4)); |
||||
int src1_data_2 = *((__global int *)((__global char *)src1 + src1_index + 8)); |
||||
|
||||
int src2_data_0 = src2.x; |
||||
int src2_data_1 = src2.y; |
||||
int src2_data_2 = src2.z; |
||||
|
||||
int data_0 = *((__global int *)((__global char *)dst + dst_index + 0)); |
||||
int data_1 = *((__global int *)((__global char *)dst + dst_index + 4)); |
||||
int data_2 = *((__global int *)((__global char *)dst + dst_index + 8)); |
||||
|
||||
int tmp_data_0 = convert_int_sat(abs_diff(src1_data_0, src2_data_0)); |
||||
int tmp_data_1 = convert_int_sat(abs_diff(src1_data_1, src2_data_1)); |
||||
int tmp_data_2 = convert_int_sat(abs_diff(src1_data_2, src2_data_2)); |
||||
|
||||
*((__global int *)((__global char *)dst + dst_index + 0))= tmp_data_0; |
||||
*((__global int *)((__global char *)dst + dst_index + 4))= tmp_data_1; |
||||
*((__global int *)((__global char *)dst + dst_index + 8))= tmp_data_2; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C3_D5 (__global float *src1, int src1_step, int src1_offset, |
||||
__global float *dst, int dst_step, int dst_offset, |
||||
float4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x * 12) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x * 12)); |
||||
|
||||
float src1_data_0 = *((__global float *)((__global char *)src1 + src1_index + 0)); |
||||
float src1_data_1 = *((__global float *)((__global char *)src1 + src1_index + 4)); |
||||
float src1_data_2 = *((__global float *)((__global char *)src1 + src1_index + 8)); |
||||
|
||||
float src2_data_0 = src2.x; |
||||
float src2_data_1 = src2.y; |
||||
float src2_data_2 = src2.z; |
||||
|
||||
float data_0 = *((__global float *)((__global char *)dst + dst_index + 0)); |
||||
float data_1 = *((__global float *)((__global char *)dst + dst_index + 4)); |
||||
float data_2 = *((__global float *)((__global char *)dst + dst_index + 8)); |
||||
|
||||
float tmp_data_0 = fabs(src1_data_0 - src2_data_0); |
||||
float tmp_data_1 = fabs(src1_data_1 - src2_data_1); |
||||
float tmp_data_2 = fabs(src1_data_2 - src2_data_2); |
||||
|
||||
*((__global float *)((__global char *)dst + dst_index + 0))= tmp_data_0; |
||||
*((__global float *)((__global char *)dst + dst_index + 4))= tmp_data_1; |
||||
*((__global float *)((__global char *)dst + dst_index + 8))= tmp_data_2; |
||||
} |
||||
} |
||||
|
||||
#if defined (DOUBLE_SUPPORT) |
||||
__kernel void arithm_s_absdiff_C3_D6 (__global double *src1, int src1_step, int src1_offset, |
||||
__global double *dst, int dst_step, int dst_offset, |
||||
double4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x * 24) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, dst_offset + (x * 24)); |
||||
|
||||
double src1_data_0 = *((__global double *)((__global char *)src1 + src1_index + 0 )); |
||||
double src1_data_1 = *((__global double *)((__global char *)src1 + src1_index + 8 )); |
||||
double src1_data_2 = *((__global double *)((__global char *)src1 + src1_index + 16)); |
||||
|
||||
double src2_data_0 = src2.x; |
||||
double src2_data_1 = src2.y; |
||||
double src2_data_2 = src2.z; |
||||
|
||||
double data_0 = *((__global double *)((__global char *)dst + dst_index + 0 )); |
||||
double data_1 = *((__global double *)((__global char *)dst + dst_index + 8 )); |
||||
double data_2 = *((__global double *)((__global char *)dst + dst_index + 16)); |
||||
|
||||
double tmp_data_0 = fabs(src1_data_0 - src2_data_0); |
||||
double tmp_data_1 = fabs(src1_data_1 - src2_data_1); |
||||
double tmp_data_2 = fabs(src1_data_2 - src2_data_2); |
||||
|
||||
*((__global double *)((__global char *)dst + dst_index + 0 ))= tmp_data_0; |
||||
*((__global double *)((__global char *)dst + dst_index + 8 ))= tmp_data_1; |
||||
*((__global double *)((__global char *)dst + dst_index + 16))= tmp_data_2; |
||||
} |
||||
} |
||||
#endif |
||||
__kernel void arithm_s_absdiff_C4_D0 (__global uchar *src1, int src1_step, int src1_offset, |
||||
__global uchar *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 2) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 2) + dst_offset); |
||||
|
||||
uchar4 src_data1 = *((__global uchar4 *)(src1 + src1_index)); |
||||
|
||||
uchar4 data = convert_uchar4_sat(abs_diff(convert_int4_sat(src_data1), src2)); |
||||
|
||||
*((__global uchar4 *)(dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C4_D2 (__global ushort *src1, int src1_step, int src1_offset, |
||||
__global ushort *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 3) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 3) + dst_offset); |
||||
|
||||
ushort4 src_data1 = *((__global ushort4 *)((__global char *)src1 + src1_index)); |
||||
|
||||
ushort4 data = convert_ushort4_sat(abs_diff(convert_int4_sat(src_data1), src2)); |
||||
|
||||
*((__global ushort4 *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C4_D3 (__global short *src1, int src1_step, int src1_offset, |
||||
__global short *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 3) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 3) + dst_offset); |
||||
|
||||
short4 src_data1 = *((__global short4 *)((__global char *)src1 + src1_index)); |
||||
|
||||
short4 data = convert_short4_sat(abs_diff(convert_int4_sat(src_data1), src2)); |
||||
|
||||
*((__global short4 *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C4_D4 (__global int *src1, int src1_step, int src1_offset, |
||||
__global int *dst, int dst_step, int dst_offset, |
||||
int4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 4) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 4) + dst_offset); |
||||
|
||||
int4 src_data1 = *((__global int4 *)((__global char *)src1 + src1_index)); |
||||
|
||||
int4 data = convert_int4_sat(abs_diff(src_data1, src2)); |
||||
|
||||
*((__global int4 *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
__kernel void arithm_s_absdiff_C4_D5 (__global float *src1, int src1_step, int src1_offset, |
||||
__global float *dst, int dst_step, int dst_offset, |
||||
float4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 4) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 4) + dst_offset); |
||||
|
||||
float4 src_data1 = *((__global float4 *)((__global char *)src1 + src1_index)); |
||||
|
||||
float4 data = fabs(src_data1 - src2); |
||||
|
||||
*((__global float4 *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
|
||||
#if defined (DOUBLE_SUPPORT) |
||||
__kernel void arithm_s_absdiff_C4_D6 (__global double *src1, int src1_step, int src1_offset, |
||||
__global double *dst, int dst_step, int dst_offset, |
||||
double4 src2, int rows, int cols, int dst_step1) |
||||
{ |
||||
|
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, (x << 5) + src1_offset); |
||||
int dst_index = mad24(y, dst_step, (x << 5) + dst_offset); |
||||
|
||||
double4 src_data1 = *((__global double4 *)((__global char *)src1 + src1_index)); |
||||
|
||||
double4 data = fabs(src_data1 - src2); |
||||
|
||||
*((__global double4 *)((__global char *)dst + dst_index)) = data; |
||||
} |
||||
} |
||||
#endif |
||||
@ -0,0 +1,79 @@ |
||||
/*M/////////////////////////////////////////////////////////////////////////////////////// |
||||
// |
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. |
||||
// |
||||
// By downloading, copying, installing or using the software you agree to this license. |
||||
// If you do not agree to this license, do not download, install, |
||||
// copy or use the software. |
||||
// |
||||
// |
||||
// License Agreement |
||||
// For Open Source Computer Vision Library |
||||
// |
||||
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. |
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. |
||||
// Third party copyrights are property of their respective owners. |
||||
// |
||||
// @Authors |
||||
// Jia Haipeng, jiahaipeng95@gmail.com |
||||
// |
||||
// Redistribution and use in source and binary forms, with or without modification, |
||||
// are permitted provided that the following conditions are met: |
||||
// |
||||
// * Redistribution's of source code must retain the above copyright notice, |
||||
// this list of conditions and the following disclaimer. |
||||
// |
||||
// * Redistribution's in binary form must reproduce the above copyright notice, |
||||
// this list of conditions and the following disclaimer in the documentation |
||||
// and/or other oclMaterials provided with the distribution. |
||||
// |
||||
// * The name of the copyright holders may not be used to endorse or promote products |
||||
// derived from this software without specific prior written permission. |
||||
// |
||||
// This software is provided by the copyright holders and contributors as is and |
||||
// any express or implied warranties, including, but not limited to, the implied |
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed. |
||||
// In no event shall the Intel Corporation or contributors be liable for any direct, |
||||
// indirect, incidental, special, exemplary, or consequential damages |
||||
// (including, but not limited to, procurement of substitute goods or services; |
||||
// loss of use, data, or profits; or business interruption) however caused |
||||
// and on any theory of liability, whether in contract, strict liability, |
||||
// or tort (including negligence or otherwise) arising in any way out of |
||||
// the use of this software, even if advised of the possibility of such damage. |
||||
// |
||||
//M*/ |
||||
|
||||
#if defined (DOUBLE_SUPPORT) |
||||
#ifdef cl_khr_fp64 |
||||
#pragma OPENCL EXTENSION cl_khr_fp64:enable |
||||
#elif defined (cl_amd_fp64) |
||||
#pragma OPENCL EXTENSION cl_amd_fp64:enable |
||||
#endif |
||||
#endif |
||||
|
||||
////////////////////////////////////////////////////////////////////////////////// |
||||
///////////////////////////////// add with mask ////////////////////////////////// |
||||
////////////////////////////////////////////////////////////////////////////////// |
||||
|
||||
__kernel void arithm_binary_op_mat_mask(__global T * src1, int src1_step, int src1_offset, |
||||
__global T * src2, int src2_step, int src2_offset, |
||||
__global uchar * mask, int mask_step, int mask_offset, |
||||
__global T * dst, int dst_step, int dst_offset, |
||||
int cols, int rows) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols && y < rows) |
||||
{ |
||||
int mask_index = mad24(y, mask_step, x + mask_offset); |
||||
if (mask[mask_index]) |
||||
{ |
||||
int src1_index = mad24(y, src1_step, x + src1_offset); |
||||
int src2_index = mad24(y, src2_step, x + src2_offset); |
||||
int dst_index = mad24(y, dst_step, dst_offset + x); |
||||
|
||||
dst[dst_index] = convertToT(convertToWT(src1[src1_index]) Operation convertToWT(src2[src2_index])); |
||||
} |
||||
} |
||||
} |
||||
@ -0,0 +1,74 @@ |
||||
/*M/////////////////////////////////////////////////////////////////////////////////////// |
||||
// |
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. |
||||
// |
||||
// By downloading, copying, installing or using the software you agree to this license. |
||||
// If you do not agree to this license, do not download, install, |
||||
// copy or use the software. |
||||
// |
||||
// |
||||
// License Agreement |
||||
// For Open Source Computer Vision Library |
||||
// |
||||
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. |
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. |
||||
// Third party copyrights are property of their respective owners. |
||||
// |
||||
// @Authors |
||||
// Jia Haipeng, jiahaipeng95@gmail.com |
||||
// |
||||
// Redistribution and use in source and binary forms, with or without modification, |
||||
// are permitted provided that the following conditions are met: |
||||
// |
||||
// * Redistribution's of source code must retain the above copyright notice, |
||||
// this list of conditions and the following disclaimer. |
||||
// |
||||
// * Redistribution's in binary form must reproduce the above copyright notice, |
||||
// this list of conditions and the following disclaimer in the documentation |
||||
// and/or other oclMaterials provided with the distribution. |
||||
// |
||||
// * The name of the copyright holders may not be used to endorse or promote products |
||||
// derived from this software without specific prior written permission. |
||||
// |
||||
// This software is provided by the copyright holders and contributors as is and |
||||
// any express or implied warranties, including, but not limited to, the implied |
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed. |
||||
// In no event shall the Intel Corporation or contributors be liable for any direct, |
||||
// indirect, incidental, special, exemplary, or consequential damages |
||||
// (including, but not limited to, procurement of substitute goods or services; |
||||
// loss of use, data, or profits; or business interruption) however caused |
||||
// and on any theory of liability, whether in contract, strict liability, |
||||
// or tort (including negligence or otherwise) arising in any way out of |
||||
// the use of this software, even if advised of the possibility of such damage. |
||||
// |
||||
//M*/ |
||||
|
||||
#if defined (DOUBLE_SUPPORT) |
||||
#ifdef cl_khr_fp64 |
||||
#pragma OPENCL EXTENSION cl_khr_fp64:enable |
||||
#elif defined (cl_amd_fp64) |
||||
#pragma OPENCL EXTENSION cl_amd_fp64:enable |
||||
#endif |
||||
#endif |
||||
|
||||
////////////////////////////////////////////////////////////////////////////////////////////////////// |
||||
/////////////////////////////////////////////addWeighted////////////////////////////////////////////// |
||||
/////////////////////////////////////////////////////////////////////////////////////////////////////// |
||||
|
||||
__kernel void arithm_compare(__global T * src1, int src1_step1, int src1_offset1, |
||||
__global T * src2, int src2_step1, int src2_offset1, |
||||
__global uchar * dst, int dst_step1, int dst_offset1, |
||||
int cols1, int rows) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (x < cols1 && y < rows) |
||||
{ |
||||
int src1_index = mad24(y, src1_step1, x + src1_offset1); |
||||
int src2_index = mad24(y, src2_step1, x + src2_offset1); |
||||
int dst_index = mad24(y, dst_step1, x + dst_offset1); |
||||
|
||||
dst[dst_index] = convert_uchar(src1[src1_index] Operation src2[src2_index] ? 255 : 0); |
||||
} |
||||
} |
||||
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in new issue