--- /dev/null
+/*#******************************************************************************
+ ** 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.
+ **
+ **
+ ** bioinspired : interfaces allowing OpenCV users to integrate Human Vision System models. Presented models originate from Jeanny Herault's original research and have been reused and adapted by the author&collaborators for computed vision applications since his thesis with Alice Caplier at Gipsa-Lab.
+ ** Use: extract still images & image sequences features, from contours details to motion spatio-temporal features, etc. for high level visual scene analysis. Also contribute to image enhancement/compression such as tone mapping.
+ **
+ ** Maintainers : Listic lab (code author current affiliation & applications) and Gipsa Lab (original research origins & applications)
+ **
+ ** Creation - enhancement process 2007-2011
+ ** Author: Alexandre Benoit (benoit.alexandre.vision@gmail.com), LISTIC lab, Annecy le vieux, France
+ **
+ ** Theses algorithm have been developped by Alexandre BENOIT since his thesis with Alice Caplier at Gipsa-Lab (www.gipsa-lab.inpg.fr) and the research he pursues at LISTIC Lab (www.listic.univ-savoie.fr).
+ ** Refer to the following research paper for more information:
+ ** Benoit A., Caplier A., Durette B., Herault, J., "USING HUMAN VISUAL SYSTEM MODELING FOR BIO-INSPIRED LOW LEVEL IMAGE PROCESSING", Elsevier, Computer Vision and Image Understanding 114 (2010), pp. 758-773, DOI: http://dx.doi.org/10.1016/j.cviu.2010.01.011
+ ** This work have been carried out thanks to Jeanny Herault who's research and great discussions are the basis of all this work, please take a look at his book:
+ ** Vision: Images, Signals and Neural Networks: Models of Neural Processing in Visual Perception (Progress in Neural Processing),By: Jeanny Herault, ISBN: 9814273686. WAPI (Tower ID): 113266891.
+ **
+ ** The retina filter includes the research contributions of phd/research collegues from which code has been redrawn by the author :
+ ** _take a look at the retinacolor.hpp module to discover Brice Chaix de Lavarene color mosaicing/demosaicing and the reference paper:
+ ** ====> B. Chaix de Lavarene, D. Alleysson, B. Durette, J. Herault (2007). "Efficient demosaicing through recursive filtering", IEEE International Conference on Image Processing ICIP 2007
+ ** _take a look at imagelogpolprojection.hpp to discover retina spatial log sampling which originates from Barthelemy Durette phd with Jeanny Herault. A Retina / V1 cortex projection is also proposed and originates from Jeanny's discussions.
+ ** ====> more informations in the above cited Jeanny Heraults's book.
+ **
+ ** License Agreement
+ ** For Open Source Computer Vision Library
+ **
+ ** Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+ ** Copyright (C) 2008-2011, Willow Garage Inc., all rights reserved.
+ **
+ ** For Human Visual System tools (bioinspired)
+ ** Copyright (C) 2007-2011, LISTIC Lab, Annecy le Vieux and GIPSA Lab, Grenoble, France, all rights reserved.
+ **
+ ** Third party copyrights are property of their respective owners.
+ **
+ ** Redistribution and use in source and binary forms, with or without modification,
+ ** are permitted provided that the following conditions are met:
+ **
+ ** * Redistributions of source code must retain the above copyright notice,
+ ** this list of conditions and the following disclaimer.
+ **
+ ** * Redistributions in binary form must reproduce the above copyright notice,
+ ** this list of conditions and the following disclaimer in the documentation
+ ** and/or other materials provided with the distribution.
+ **
+ ** * The name of the copyright holders may not be used to endorse or promote products
+ ** derived from this software without specific prior written permission.
+ **
+ ** This software is provided by the copyright holders and contributors "as is" and
+ ** any express or implied warranties, including, but not limited to, the implied
+ ** warranties of merchantability and fitness for a particular purpose are disclaimed.
+ ** In no event shall the Intel Corporation or contributors be liable for any direct,
+ ** indirect, incidental, special, exemplary, or consequential damages
+ ** (including, but not limited to, procurement of substitute goods or services;
+ ** loss of use, data, or profits; or business interruption) however caused
+ ** and on any theory of liability, whether in contract, strict liability,
+ ** or tort (including negligence or otherwise) arising in any way out of
+ ** the use of this software, even if advised of the possibility of such damage.
+ *******************************************************************************/
+
+/*
+ * Retina.cpp
+ *
+ * Created on: Jul 19, 2011
+ * Author: Alexandre Benoit
+ */
+#include "precomp.hpp"
+#include "retinafilter.hpp"
+#include <cstdio>
+#include <sstream>
+#include <valarray>
+
+namespace cv
+{
+namespace bioinspired
+{
+
+class RetinaImpl : public Retina
+{
+public:
+ /**
+ * Main constructor with most commun use setup : create an instance of color ready retina model
+ * @param inputSize : the input frame size
+ */
+ RetinaImpl(Size inputSize);
+
+ /**
+ * Complete Retina filter constructor which allows all basic structural parameters definition
+ * @param inputSize : the input frame size
+ * @param colorMode : the chosen processing mode : with or without color processing
+ * @param colorSamplingMethod: specifies which kind of color sampling will be used
+ * @param useRetinaLogSampling: activate retina log sampling, if true, the 2 following parameters can be used
+ * @param reductionFactor: only usefull if param useRetinaLogSampling=true, specifies the reduction factor of the output frame (as the center (fovea) is high resolution and corners can be underscaled, then a reduction of the output is allowed without precision leak
+ * @param samplingStrenght: only usefull if param useRetinaLogSampling=true, specifies the strenght of the log scale that is applied
+ */
+ RetinaImpl(Size inputSize, const bool colorMode, int colorSamplingMethod=RETINA_COLOR_BAYER, const bool useRetinaLogSampling=false, const double reductionFactor=1.0, const double samplingStrenght=10.0);
+
+ virtual ~RetinaImpl();
+ /**
+ * retreive retina input buffer size
+ */
+ Size getInputSize();
+
+ /**
+ * retreive retina output buffer size
+ */
+ Size getOutputSize();
+
+ /**
+ * try to open an XML retina parameters file to adjust current retina instance setup
+ * => if the xml file does not exist, then default setup is applied
+ * => warning, Exceptions are thrown if read XML file is not valid
+ * @param retinaParameterFile : the parameters filename
+ * @param applyDefaultSetupOnFailure : set to true if an error must be thrown on error
+ */
+ void setup(String retinaParameterFile="", const bool applyDefaultSetupOnFailure=true);
+
+
+ /**
+ * try to open an XML retina parameters file to adjust current retina instance setup
+ * => if the xml file does not exist, then default setup is applied
+ * => warning, Exceptions are thrown if read XML file is not valid
+ * @param fs : the open Filestorage which contains retina parameters
+ * @param applyDefaultSetupOnFailure : set to true if an error must be thrown on error
+ */
+ void setup(cv::FileStorage &fs, const bool applyDefaultSetupOnFailure=true);
+
+ /**
+ * try to open an XML retina parameters file to adjust current retina instance setup
+ * => if the xml file does not exist, then default setup is applied
+ * => warning, Exceptions are thrown if read XML file is not valid
+ * @param newParameters : a parameters structures updated with the new target configuration
+ * @param applyDefaultSetupOnFailure : set to true if an error must be thrown on error
+ */
+ void setup(Retina::RetinaParameters newParameters);
+
+ /**
+ * @return the current parameters setup
+ */
+ struct Retina::RetinaParameters getParameters();
+
+ /**
+ * parameters setup display method
+ * @return a string which contains formatted parameters information
+ */
+ const String printSetup();
+
+ /**
+ * write xml/yml formated parameters information
+ * @rparam fs : the filename of the xml file that will be open and writen with formatted parameters information
+ */
+ virtual void write( String fs ) const;
+
+
+ /**
+ * write xml/yml formated parameters information
+ * @param fs : a cv::Filestorage object ready to be filled
+ */
+ virtual void write( FileStorage& fs ) const;
+
+ /**
+ * setup the OPL and IPL parvo channels (see biologocal model)
+ * OPL is referred as Outer Plexiform Layer of the retina, it allows the spatio-temporal filtering which withens the spectrum and reduces spatio-temporal noise while attenuating global luminance (low frequency energy)
+ * IPL parvo is the OPL next processing stage, it refers to Inner Plexiform layer of the retina, it allows high contours sensitivity in foveal vision.
+ * for more informations, please have a look at the paper Benoit A., Caplier A., Durette B., Herault, J., "USING HUMAN VISUAL SYSTEM MODELING FOR BIO-INSPIRED LOW LEVEL IMAGE PROCESSING", Elsevier, Computer Vision and Image Understanding 114 (2010), pp. 758-773, DOI: http://dx.doi.org/10.1016/j.cviu.2010.01.011
+ * @param colorMode : specifies if (true) color is processed of not (false) to then processing gray level image
+ * @param normaliseOutput : specifies if (true) output is rescaled between 0 and 255 of not (false)
+ * @param photoreceptorsLocalAdaptationSensitivity: the photoreceptors sensitivity renage is 0-1 (more log compression effect when value increases)
+ * @param photoreceptorsTemporalConstant: the time constant of the first order low pass filter of the photoreceptors, use it to cut high temporal frequencies (noise or fast motion), unit is frames, typical value is 1 frame
+ * @param photoreceptorsSpatialConstant: the spatial constant of the first order low pass filter of the photoreceptors, use it to cut high spatial frequencies (noise or thick contours), unit is pixels, typical value is 1 pixel
+ * @param horizontalCellsGain: gain of the horizontal cells network, if 0, then the mean value of the output is zero, if the parameter is near 1, then, the luminance is not filtered and is still reachable at the output, typicall value is 0
+ * @param HcellsTemporalConstant: the time constant of the first order low pass filter of the horizontal cells, use it to cut low temporal frequencies (local luminance variations), unit is frames, typical value is 1 frame, as the photoreceptors
+ * @param HcellsSpatialConstant: the spatial constant of the first order low pass filter of the horizontal cells, use it to cut low spatial frequencies (local luminance), unit is pixels, typical value is 5 pixel, this value is also used for local contrast computing when computing the local contrast adaptation at the ganglion cells level (Inner Plexiform Layer parvocellular channel model)
+ * @param ganglionCellsSensitivity: the compression strengh of the ganglion cells local adaptation output, set a value between 160 and 250 for best results, a high value increases more the low value sensitivity... and the output saturates faster, recommended value: 230
+ */
+ void setupOPLandIPLParvoChannel(const bool colorMode=true, const bool normaliseOutput = true, const float photoreceptorsLocalAdaptationSensitivity=0.7, const float photoreceptorsTemporalConstant=0.5, const float photoreceptorsSpatialConstant=0.53, const float horizontalCellsGain=0, const float HcellsTemporalConstant=1, const float HcellsSpatialConstant=7, const float ganglionCellsSensitivity=0.7);
+
+ /**
+ * set parameters values for the Inner Plexiform Layer (IPL) magnocellular channel
+ * this channel processes signals outpint from OPL processing stage in peripheral vision, it allows motion information enhancement. It is decorrelated from the details channel. See reference paper for more details.
+ * @param normaliseOutput : specifies if (true) output is rescaled between 0 and 255 of not (false)
+ * @param parasolCells_beta: the low pass filter gain used for local contrast adaptation at the IPL level of the retina (for ganglion cells local adaptation), typical value is 0
+ * @param parasolCells_tau: the low pass filter time constant used for local contrast adaptation at the IPL level of the retina (for ganglion cells local adaptation), unit is frame, typical value is 0 (immediate response)
+ * @param parasolCells_k: the low pass filter spatial constant used for local contrast adaptation at the IPL level of the retina (for ganglion cells local adaptation), unit is pixels, typical value is 5
+ * @param amacrinCellsTemporalCutFrequency: the time constant of the first order high pass fiter of the magnocellular way (motion information channel), unit is frames, tipicall value is 5
+ * @param V0CompressionParameter: the compression strengh of the ganglion cells local adaptation output, set a value between 160 and 250 for best results, a high value increases more the low value sensitivity... and the output saturates faster, recommended value: 200
+ * @param localAdaptintegration_tau: specifies the temporal constant of the low pas filter involved in the computation of the local "motion mean" for the local adaptation computation
+ * @param localAdaptintegration_k: specifies the spatial constant of the low pas filter involved in the computation of the local "motion mean" for the local adaptation computation
+ */
+ void setupIPLMagnoChannel(const bool normaliseOutput = true, const float parasolCells_beta=0, const float parasolCells_tau=0, const float parasolCells_k=7, const float amacrinCellsTemporalCutFrequency=1.2, const float V0CompressionParameter=0.95, const float localAdaptintegration_tau=0, const float localAdaptintegration_k=7);
+
+ /**
+ * method which allows retina to be applied on an input image, after run, encapsulated retina module is ready to deliver its outputs using dedicated acccessors, see getParvo and getMagno methods
+ * @param inputImage : the input cv::Mat image to be processed, can be gray level or BGR coded in any format (from 8bit to 16bits)
+ */
+ void run(InputArray inputImage);
+
+ /**
+ * method that applies a luminance correction (initially High Dynamic Range (HDR) tone mapping) using only the 2 local adaptation stages of the retina parvo channel : photoreceptors level and ganlion cells level. Spatio temporal filtering is applied but limited to temporal smoothing and eventually high frequencies attenuation. This is a lighter method than the one available using the regular run method. It is then faster but it does not include complete temporal filtering nor retina spectral whitening. This is an adptation of the original still image HDR tone mapping algorithm of David Alleyson, Sabine Susstruck and Laurence Meylan's work, please cite:
+ * -> Meylan L., Alleysson D., and Susstrunk S., A Model of Retinal Local Adaptation for the Tone Mapping of Color Filter Array Images, Journal of Optical Society of America, A, Vol. 24, N 9, September, 1st, 2007, pp. 2807-2816
+ @param inputImage the input image to process RGB or gray levels
+ @param outputToneMappedImage the output tone mapped image
+ */
+ void applyFastToneMapping(InputArray inputImage, OutputArray outputToneMappedImage);
+
+ /**
+ * accessor of the details channel of the retina (models foveal vision)
+ * @param retinaOutput_parvo : the output buffer (reallocated if necessary), this output is rescaled for standard 8bits image processing use in OpenCV
+ */
+ void getParvo(OutputArray retinaOutput_parvo);
+
+ /**
+ * accessor of the details channel of the retina (models foveal vision)
+ * @param retinaOutput_parvo : a cv::Mat header filled with the internal parvo buffer of the retina module. This output is the original retina filter model output, without any quantification or rescaling
+ */
+ void getParvoRAW(OutputArray retinaOutput_parvo);
+
+ /**
+ * accessor of the motion channel of the retina (models peripheral vision)
+ * @param retinaOutput_magno : the output buffer (reallocated if necessary), this output is rescaled for standard 8bits image processing use in OpenCV
+ */
+ void getMagno(OutputArray retinaOutput_magno);
+
+ /**
+ * accessor of the motion channel of the retina (models peripheral vision)
+ * @param retinaOutput_magno : a cv::Mat header filled with the internal retina magno buffer of the retina module. This output is the original retina filter model output, without any quantification or rescaling
+ */
+ void getMagnoRAW(OutputArray retinaOutput_magno);
+
+ // original API level data accessors : get buffers addresses from a Mat header, similar to getParvoRAW and getMagnoRAW...
+ const Mat getMagnoRAW() const;
+ const Mat getParvoRAW() const;
+
+ /**
+ * activate color saturation as the final step of the color demultiplexing process
+ * -> this saturation is a sigmoide function applied to each channel of the demultiplexed image.
+ * @param saturateColors: boolean that activates color saturation (if true) or desactivate (if false)
+ * @param colorSaturationValue: the saturation factor
+ */
+ void setColorSaturation(const bool saturateColors=true, const float colorSaturationValue=4.0);
+
+ /**
+ * clear all retina buffers (equivalent to opening the eyes after a long period of eye close ;o)
+ */
+ void clearBuffers();
+
+ /**
+ * Activate/desactivate the Magnocellular pathway processing (motion information extraction), by default, it is activated
+ * @param activate: true if Magnocellular output should be activated, false if not
+ */
+ void activateMovingContoursProcessing(const bool activate);
+
+ /**
+ * Activate/desactivate the Parvocellular pathway processing (contours information extraction), by default, it is activated
+ * @param activate: true if Parvocellular (contours information extraction) output should be activated, false if not
+ */
+ void activateContoursProcessing(const bool activate);
+private:
+
+ // Parameteres setup members
+ RetinaParameters _retinaParameters; // structure of parameters
+
+ // Retina model related modules
+ std::valarray<float> _inputBuffer; //!< buffer used to convert input cv::Mat to internal retina buffers format (valarrays)
+
+ // pointer to retina model
+ RetinaFilter* _retinaFilter; //!< the pointer to the retina module, allocated with instance construction
+
+ //! private method called by constructors, gathers their parameters and use them in a unified way
+ void _init(const Size inputSize, const bool colorMode, int colorSamplingMethod=RETINA_COLOR_BAYER, const bool useRetinaLogSampling=false, const double reductionFactor=1.0, const double samplingStrenght=10.0);
+
+ /**
+ * exports a valarray buffer outing from bioinspired objects to a cv::Mat in CV_8UC1 (gray level picture) or CV_8UC3 (color) format
+ * @param grayMatrixToConvert the valarray to export to OpenCV
+ * @param nbRows : the number of rows of the valarray flatten matrix
+ * @param nbColumns : the number of rows of the valarray flatten matrix
+ * @param colorMode : a flag which mentions if matrix is color (true) or graylevel (false)
+ * @param outBuffer : the output matrix which is reallocated to satisfy Retina output buffer dimensions
+ */
+ void _convertValarrayBuffer2cvMat(const std::valarray<float> &grayMatrixToConvert, const unsigned int nbRows, const unsigned int nbColumns, const bool colorMode, OutputArray outBuffer);
+
+ /**
+ * convert a cv::Mat to a valarray buffer in float format
+ * @param inputMatToConvert : the OpenCV cv::Mat that has to be converted to gray or RGB valarray buffer that will be processed by the retina model
+ * @param outputValarrayMatrix : the output valarray
+ * @return the input image color mode (color=true, gray levels=false)
+ */
+ bool _convertCvMat2ValarrayBuffer(InputArray inputMatToConvert, std::valarray<float> &outputValarrayMatrix);
+
+
+};
+
+// smart pointers allocation :
+Ptr<Retina> createRetina(Size inputSize){ return new RetinaImpl(inputSize); }
+Ptr<Retina> createRetina(Size inputSize, const bool colorMode, int colorSamplingMethod, const bool useRetinaLogSampling, const double reductionFactor, const double samplingStrenght){return new RetinaImpl(inputSize, colorMode, colorSamplingMethod, useRetinaLogSampling, reductionFactor, samplingStrenght);}
+
+
+// RetinaImpl code
+RetinaImpl::RetinaImpl(const cv::Size inputSz)
+{
+ _retinaFilter = 0;
+ _init(inputSz, true, RETINA_COLOR_BAYER, false);
+}
+
+RetinaImpl::RetinaImpl(const cv::Size inputSz, const bool colorMode, int colorSamplingMethod, const bool useRetinaLogSampling, const double reductionFactor, const double samplingStrenght)
+{
+ _retinaFilter = 0;
+ _init(inputSz, colorMode, colorSamplingMethod, useRetinaLogSampling, reductionFactor, samplingStrenght);
+};
+
+RetinaImpl::~RetinaImpl()
+{
+ if (_retinaFilter)
+ delete _retinaFilter;
+}
+
+/**
+* retreive retina input buffer size
+*/
+Size RetinaImpl::getInputSize(){return cv::Size(_retinaFilter->getInputNBcolumns(), _retinaFilter->getInputNBrows());}
+
+/**
+* retreive retina output buffer size
+*/
+Size RetinaImpl::getOutputSize(){return cv::Size(_retinaFilter->getOutputNBcolumns(), _retinaFilter->getOutputNBrows());}
+
+
+void RetinaImpl::setColorSaturation(const bool saturateColors, const float colorSaturationValue)
+{
+ _retinaFilter->setColorSaturation(saturateColors, colorSaturationValue);
+}
+
+struct Retina::RetinaParameters RetinaImpl::getParameters(){return _retinaParameters;}
+
+void RetinaImpl::setup(String retinaParameterFile, const bool applyDefaultSetupOnFailure)
+{
+ try
+ {
+ // opening retinaParameterFile in read mode
+ cv::FileStorage fs(retinaParameterFile, cv::FileStorage::READ);
+ setup(fs, applyDefaultSetupOnFailure);
+ }
+ catch(Exception &e)
+ {
+ printf("Retina::setup: wrong/unappropriate xml parameter file : error report :`n=>%s\n", e.what());
+ if (applyDefaultSetupOnFailure)
+ {
+ printf("Retina::setup: resetting retina with default parameters\n");
+ setupOPLandIPLParvoChannel();
+ setupIPLMagnoChannel();
+ }
+ else
+ {
+ printf("=> keeping current parameters\n");
+ }
+ }
+}
+
+void RetinaImpl::setup(cv::FileStorage &fs, const bool applyDefaultSetupOnFailure)
+{
+ try
+ {
+ // read parameters file if it exists or apply default setup if asked for
+ if (!fs.isOpened())
+ {
+ printf("Retina::setup: provided parameters file could not be open... skeeping configuration\n");
+ return;
+ // implicit else case : retinaParameterFile could be open (it exists at least)
+ }
+ // OPL and Parvo init first... update at the same time the parameters structure and the retina core
+ cv::FileNode rootFn = fs.root(), currFn=rootFn["OPLandIPLparvo"];
+ currFn["colorMode"]>>_retinaParameters.OPLandIplParvo.colorMode;
+ currFn["normaliseOutput"]>>_retinaParameters.OPLandIplParvo.normaliseOutput;
+ currFn["photoreceptorsLocalAdaptationSensitivity"]>>_retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity;
+ currFn["photoreceptorsTemporalConstant"]>>_retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant;
+ currFn["photoreceptorsSpatialConstant"]>>_retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant;
+ currFn["horizontalCellsGain"]>>_retinaParameters.OPLandIplParvo.horizontalCellsGain;
+ currFn["hcellsTemporalConstant"]>>_retinaParameters.OPLandIplParvo.hcellsTemporalConstant;
+ currFn["hcellsSpatialConstant"]>>_retinaParameters.OPLandIplParvo.hcellsSpatialConstant;
+ currFn["ganglionCellsSensitivity"]>>_retinaParameters.OPLandIplParvo.ganglionCellsSensitivity;
+ setupOPLandIPLParvoChannel(_retinaParameters.OPLandIplParvo.colorMode, _retinaParameters.OPLandIplParvo.normaliseOutput, _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity, _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant, _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant, _retinaParameters.OPLandIplParvo.horizontalCellsGain, _retinaParameters.OPLandIplParvo.hcellsTemporalConstant, _retinaParameters.OPLandIplParvo.hcellsSpatialConstant, _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity);
+
+ // init retina IPL magno setup... update at the same time the parameters structure and the retina core
+ currFn=rootFn["IPLmagno"];
+ currFn["normaliseOutput"]>>_retinaParameters.IplMagno.normaliseOutput;
+ currFn["parasolCells_beta"]>>_retinaParameters.IplMagno.parasolCells_beta;
+ currFn["parasolCells_tau"]>>_retinaParameters.IplMagno.parasolCells_tau;
+ currFn["parasolCells_k"]>>_retinaParameters.IplMagno.parasolCells_k;
+ currFn["amacrinCellsTemporalCutFrequency"]>>_retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency;
+ currFn["V0CompressionParameter"]>>_retinaParameters.IplMagno.V0CompressionParameter;
+ currFn["localAdaptintegration_tau"]>>_retinaParameters.IplMagno.localAdaptintegration_tau;
+ currFn["localAdaptintegration_k"]>>_retinaParameters.IplMagno.localAdaptintegration_k;
+
+ setupIPLMagnoChannel(_retinaParameters.IplMagno.normaliseOutput, _retinaParameters.IplMagno.parasolCells_beta, _retinaParameters.IplMagno.parasolCells_tau, _retinaParameters.IplMagno.parasolCells_k, _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency,_retinaParameters.IplMagno.V0CompressionParameter, _retinaParameters.IplMagno.localAdaptintegration_tau, _retinaParameters.IplMagno.localAdaptintegration_k);
+
+ }catch(Exception &e)
+ {
+ printf("RetinaImpl::setup: resetting retina with default parameters\n");
+ if (applyDefaultSetupOnFailure)
+ {
+ setupOPLandIPLParvoChannel();
+ setupIPLMagnoChannel();
+ }
+ printf("Retina::setup: wrong/unappropriate xml parameter file : error report :`n=>%s\n", e.what());
+ printf("=> keeping current parameters\n");
+ }
+
+ // report current configuration
+ printf("%s\n", printSetup().c_str());
+}
+
+void RetinaImpl::setup(Retina::RetinaParameters newConfiguration)
+{
+ // simply copy structures
+ memcpy(&_retinaParameters, &newConfiguration, sizeof(Retina::RetinaParameters));
+ // apply setup
+ setupOPLandIPLParvoChannel(_retinaParameters.OPLandIplParvo.colorMode, _retinaParameters.OPLandIplParvo.normaliseOutput, _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity, _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant, _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant, _retinaParameters.OPLandIplParvo.horizontalCellsGain, _retinaParameters.OPLandIplParvo.hcellsTemporalConstant, _retinaParameters.OPLandIplParvo.hcellsSpatialConstant, _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity);
+ setupIPLMagnoChannel(_retinaParameters.IplMagno.normaliseOutput, _retinaParameters.IplMagno.parasolCells_beta, _retinaParameters.IplMagno.parasolCells_tau, _retinaParameters.IplMagno.parasolCells_k, _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency,_retinaParameters.IplMagno.V0CompressionParameter, _retinaParameters.IplMagno.localAdaptintegration_tau, _retinaParameters.IplMagno.localAdaptintegration_k);
+
+}
+
+const String RetinaImpl::printSetup()
+{
+ std::stringstream outmessage;
+
+ // displaying OPL and IPL parvo setup
+ outmessage<<"Current Retina instance setup :"
+ <<"\nOPLandIPLparvo"<<"{"
+ << "\n\t colorMode : " << _retinaParameters.OPLandIplParvo.colorMode
+ << "\n\t normalizeParvoOutput :" << _retinaParameters.OPLandIplParvo.normaliseOutput
+ << "\n\t photoreceptorsLocalAdaptationSensitivity : " << _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity
+ << "\n\t photoreceptorsTemporalConstant : " << _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant
+ << "\n\t photoreceptorsSpatialConstant : " << _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant
+ << "\n\t horizontalCellsGain : " << _retinaParameters.OPLandIplParvo.horizontalCellsGain
+ << "\n\t hcellsTemporalConstant : " << _retinaParameters.OPLandIplParvo.hcellsTemporalConstant
+ << "\n\t hcellsSpatialConstant : " << _retinaParameters.OPLandIplParvo.hcellsSpatialConstant
+ << "\n\t parvoGanglionCellsSensitivity : " << _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity
+ <<"}\n";
+
+ // displaying IPL magno setup
+ outmessage<<"Current Retina instance setup :"
+ <<"\nIPLmagno"<<"{"
+ << "\n\t normaliseOutput : " << _retinaParameters.IplMagno.normaliseOutput
+ << "\n\t parasolCells_beta : " << _retinaParameters.IplMagno.parasolCells_beta
+ << "\n\t parasolCells_tau : " << _retinaParameters.IplMagno.parasolCells_tau
+ << "\n\t parasolCells_k : " << _retinaParameters.IplMagno.parasolCells_k
+ << "\n\t amacrinCellsTemporalCutFrequency : " << _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency
+ << "\n\t V0CompressionParameter : " << _retinaParameters.IplMagno.V0CompressionParameter
+ << "\n\t localAdaptintegration_tau : " << _retinaParameters.IplMagno.localAdaptintegration_tau
+ << "\n\t localAdaptintegration_k : " << _retinaParameters.IplMagno.localAdaptintegration_k
+ <<"}";
+ return outmessage.str().c_str();
+}
+
+void RetinaImpl::write( String fs ) const
+{
+ FileStorage parametersSaveFile(fs, cv::FileStorage::WRITE );
+ write(parametersSaveFile);
+}
+
+void RetinaImpl::write( FileStorage& fs ) const
+{
+ if (!fs.isOpened())
+ return; // basic error case
+ fs<<"OPLandIPLparvo"<<"{";
+ fs << "colorMode" << _retinaParameters.OPLandIplParvo.colorMode;
+ fs << "normaliseOutput" << _retinaParameters.OPLandIplParvo.normaliseOutput;
+ fs << "photoreceptorsLocalAdaptationSensitivity" << _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity;
+ fs << "photoreceptorsTemporalConstant" << _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant;
+ fs << "photoreceptorsSpatialConstant" << _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant;
+ fs << "horizontalCellsGain" << _retinaParameters.OPLandIplParvo.horizontalCellsGain;
+ fs << "hcellsTemporalConstant" << _retinaParameters.OPLandIplParvo.hcellsTemporalConstant;
+ fs << "hcellsSpatialConstant" << _retinaParameters.OPLandIplParvo.hcellsSpatialConstant;
+ fs << "ganglionCellsSensitivity" << _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity;
+ fs << "}";
+ fs<<"IPLmagno"<<"{";
+ fs << "normaliseOutput" << _retinaParameters.IplMagno.normaliseOutput;
+ fs << "parasolCells_beta" << _retinaParameters.IplMagno.parasolCells_beta;
+ fs << "parasolCells_tau" << _retinaParameters.IplMagno.parasolCells_tau;
+ fs << "parasolCells_k" << _retinaParameters.IplMagno.parasolCells_k;
+ fs << "amacrinCellsTemporalCutFrequency" << _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency;
+ fs << "V0CompressionParameter" << _retinaParameters.IplMagno.V0CompressionParameter;
+ fs << "localAdaptintegration_tau" << _retinaParameters.IplMagno.localAdaptintegration_tau;
+ fs << "localAdaptintegration_k" << _retinaParameters.IplMagno.localAdaptintegration_k;
+ fs<<"}";
+}
+
+void RetinaImpl::setupOPLandIPLParvoChannel(const bool colorMode, const bool normaliseOutput, const float photoreceptorsLocalAdaptationSensitivity, const float photoreceptorsTemporalConstant, const float photoreceptorsSpatialConstant, const float horizontalCellsGain, const float HcellsTemporalConstant, const float HcellsSpatialConstant, const float ganglionCellsSensitivity)
+{
+ // retina core parameters setup
+ _retinaFilter->setColorMode(colorMode);
+ _retinaFilter->setPhotoreceptorsLocalAdaptationSensitivity(photoreceptorsLocalAdaptationSensitivity);
+ _retinaFilter->setOPLandParvoParameters(0, photoreceptorsTemporalConstant, photoreceptorsSpatialConstant, horizontalCellsGain, HcellsTemporalConstant, HcellsSpatialConstant, ganglionCellsSensitivity);
+ _retinaFilter->setParvoGanglionCellsLocalAdaptationSensitivity(ganglionCellsSensitivity);
+ _retinaFilter->activateNormalizeParvoOutput_0_maxOutputValue(normaliseOutput);
+
+ // update parameters struture
+
+ _retinaParameters.OPLandIplParvo.colorMode = colorMode;
+ _retinaParameters.OPLandIplParvo.normaliseOutput = normaliseOutput;
+ _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity = photoreceptorsLocalAdaptationSensitivity;
+ _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant = photoreceptorsTemporalConstant;
+ _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant = photoreceptorsSpatialConstant;
+ _retinaParameters.OPLandIplParvo.horizontalCellsGain = horizontalCellsGain;
+ _retinaParameters.OPLandIplParvo.hcellsTemporalConstant = HcellsTemporalConstant;
+ _retinaParameters.OPLandIplParvo.hcellsSpatialConstant = HcellsSpatialConstant;
+ _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity = ganglionCellsSensitivity;
+
+}
+
+void RetinaImpl::setupIPLMagnoChannel(const bool normaliseOutput, const float parasolCells_beta, const float parasolCells_tau, const float parasolCells_k, const float amacrinCellsTemporalCutFrequency, const float V0CompressionParameter, const float localAdaptintegration_tau, const float localAdaptintegration_k)
+{
+
+ _retinaFilter->setMagnoCoefficientsTable(parasolCells_beta, parasolCells_tau, parasolCells_k, amacrinCellsTemporalCutFrequency, V0CompressionParameter, localAdaptintegration_tau, localAdaptintegration_k);
+ _retinaFilter->activateNormalizeMagnoOutput_0_maxOutputValue(normaliseOutput);
+
+ // update parameters struture
+ _retinaParameters.IplMagno.normaliseOutput = normaliseOutput;
+ _retinaParameters.IplMagno.parasolCells_beta = parasolCells_beta;
+ _retinaParameters.IplMagno.parasolCells_tau = parasolCells_tau;
+ _retinaParameters.IplMagno.parasolCells_k = parasolCells_k;
+ _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency = amacrinCellsTemporalCutFrequency;
+ _retinaParameters.IplMagno.V0CompressionParameter = V0CompressionParameter;
+ _retinaParameters.IplMagno.localAdaptintegration_tau = localAdaptintegration_tau;
+ _retinaParameters.IplMagno.localAdaptintegration_k = localAdaptintegration_k;
+}
+
+void RetinaImpl::run(InputArray inputMatToConvert)
+{
+ // first convert input image to the compatible format : std::valarray<float>
+ const bool colorMode = _convertCvMat2ValarrayBuffer(inputMatToConvert.getMat(), _inputBuffer);
+ // process the retina
+ if (!_retinaFilter->runFilter(_inputBuffer, colorMode, false, _retinaParameters.OPLandIplParvo.colorMode && colorMode, false))
+ throw cv::Exception(-1, "RetinaImpl cannot be applied, wrong input buffer size", "RetinaImpl::run", "RetinaImpl.h", 0);
+}
+
+void RetinaImpl::applyFastToneMapping(InputArray inputImage, OutputArray outputToneMappedImage)
+{
+ // first convert input image to the compatible format :
+ const bool colorMode = _convertCvMat2ValarrayBuffer(inputImage.getMat(), _inputBuffer);
+ const unsigned int nbPixels=_retinaFilter->getOutputNBrows()*_retinaFilter->getOutputNBcolumns();
+
+ // process tone mapping
+ if (colorMode)
+ {
+ std::valarray<float> imageOutput(nbPixels*3);
+ _retinaFilter->runRGBToneMapping(_inputBuffer, imageOutput, true, _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity, _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity);
+ _convertValarrayBuffer2cvMat(imageOutput, _retinaFilter->getOutputNBrows(), _retinaFilter->getOutputNBcolumns(), true, outputToneMappedImage);
+ }else
+ {
+ std::valarray<float> imageOutput(nbPixels);
+ _retinaFilter->runGrayToneMapping(_inputBuffer, imageOutput, _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity, _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity);
+ _convertValarrayBuffer2cvMat(imageOutput, _retinaFilter->getOutputNBrows(), _retinaFilter->getOutputNBcolumns(), false, outputToneMappedImage);
+ }
+
+}
+
+void RetinaImpl::getParvo(OutputArray retinaOutput_parvo)
+{
+ if (_retinaFilter->getColorMode())
+ {
+ // reallocate output buffer (if necessary)
+ _convertValarrayBuffer2cvMat(_retinaFilter->getColorOutput(), _retinaFilter->getOutputNBrows(), _retinaFilter->getOutputNBcolumns(), true, retinaOutput_parvo);
+ }else
+ {
+ // reallocate output buffer (if necessary)
+ _convertValarrayBuffer2cvMat(_retinaFilter->getContours(), _retinaFilter->getOutputNBrows(), _retinaFilter->getOutputNBcolumns(), false, retinaOutput_parvo);
+ }
+ //retinaOutput_parvo/=255.0;
+}
+void RetinaImpl::getMagno(OutputArray retinaOutput_magno)
+{
+ // reallocate output buffer (if necessary)
+ _convertValarrayBuffer2cvMat(_retinaFilter->getMovingContours(), _retinaFilter->getOutputNBrows(), _retinaFilter->getOutputNBcolumns(), false, retinaOutput_magno);
+ //retinaOutput_magno/=255.0;
+}
+
+// original API level data accessors : copy buffers if size matches, reallocate if required
+void RetinaImpl::getMagnoRAW(OutputArray magnoOutputBufferCopy){
+ // get magno channel header
+ const cv::Mat magnoChannel=cv::Mat(getMagnoRAW());
+ // copy data
+ magnoChannel.copyTo(magnoOutputBufferCopy);
+}
+
+void RetinaImpl::getParvoRAW(OutputArray parvoOutputBufferCopy){
+ // get parvo channel header
+ const cv::Mat parvoChannel=cv::Mat(getMagnoRAW());
+ // copy data
+ parvoChannel.copyTo(parvoOutputBufferCopy);
+}
+
+// original API level data accessors : get buffers addresses...
+const Mat RetinaImpl::getMagnoRAW() const {
+ // create a cv::Mat header for the valarray
+ return Mat(_retinaFilter->getMovingContours().size(),1, CV_32F, (void*)get_data(_retinaFilter->getMovingContours()));
+
+}
+
+const Mat RetinaImpl::getParvoRAW() const {
+ if (_retinaFilter->getColorMode()) // check if color mode is enabled
+ {
+ // create a cv::Mat table (for RGB planes as a single vector)
+ return Mat(_retinaFilter->getColorOutput().size(), 1, CV_32F, (void*)get_data(_retinaFilter->getColorOutput()));
+ }
+ // otherwise, output is gray level
+ // create a cv::Mat header for the valarray
+ return Mat( _retinaFilter->getContours().size(), 1, CV_32F, (void*)get_data(_retinaFilter->getContours()));
+}
+
+// private method called by constructirs
+void RetinaImpl::_init(const cv::Size inputSz, const bool colorMode, int colorSamplingMethod, const bool useRetinaLogSampling, const double reductionFactor, const double samplingStrenght)
+{
+ // basic error check
+ if (inputSz.height*inputSz.width <= 0)
+ throw cv::Exception(-1, "Bad retina size setup : size height and with must be superior to zero", "RetinaImpl::setup", "Retina.cpp", 0);
+
+ unsigned int nbPixels=inputSz.height*inputSz.width;
+ // resize buffers if size does not match
+ _inputBuffer.resize(nbPixels*3); // buffer supports gray images but also 3 channels color buffers... (larger is better...)
+
+ // allocate the retina model
+ if (_retinaFilter)
+ delete _retinaFilter;
+ _retinaFilter = new RetinaFilter(inputSz.height, inputSz.width, colorMode, colorSamplingMethod, useRetinaLogSampling, reductionFactor, samplingStrenght);
+
++ _retinaParameters.OPLandIplParvo.colorMode = colorMode;
+ // prepare the default parameter XML file with default setup
+ setup(_retinaParameters);
+
+ // init retina
+ _retinaFilter->clearAllBuffers();
+
+ // report current configuration
+ printf("%s\n", printSetup().c_str());
+}
+
+void RetinaImpl::_convertValarrayBuffer2cvMat(const std::valarray<float> &grayMatrixToConvert, const unsigned int nbRows, const unsigned int nbColumns, const bool colorMode, OutputArray outBuffer)
+{
+ // fill output buffer with the valarray buffer
+ const float *valarrayPTR=get_data(grayMatrixToConvert);
+ if (!colorMode)
+ {
+ outBuffer.create(cv::Size(nbColumns, nbRows), CV_8U);
+ Mat outMat = outBuffer.getMat();
+ for (unsigned int i=0;i<nbRows;++i)
+ {
+ for (unsigned int j=0;j<nbColumns;++j)
+ {
+ cv::Point2d pixel(j,i);
+ outMat.at<unsigned char>(pixel)=(unsigned char)*(valarrayPTR++);
+ }
+ }
+ }else
+ {
+ const unsigned int nbPixels=nbColumns*nbRows;
+ const unsigned int doubleNBpixels=nbColumns*nbRows*2;
+ outBuffer.create(cv::Size(nbColumns, nbRows), CV_8UC3);
+ Mat outMat = outBuffer.getMat();
+ for (unsigned int i=0;i<nbRows;++i)
+ {
+ for (unsigned int j=0;j<nbColumns;++j,++valarrayPTR)
+ {
+ cv::Point2d pixel(j,i);
+ cv::Vec3b pixelValues;
+ pixelValues[2]=(unsigned char)*(valarrayPTR);
+ pixelValues[1]=(unsigned char)*(valarrayPTR+nbPixels);
+ pixelValues[0]=(unsigned char)*(valarrayPTR+doubleNBpixels);
+
+ outMat.at<cv::Vec3b>(pixel)=pixelValues;
+ }
+ }
+ }
+}
+
+bool RetinaImpl::_convertCvMat2ValarrayBuffer(InputArray inputMat, std::valarray<float> &outputValarrayMatrix)
+{
+ const Mat inputMatToConvert=inputMat.getMat();
+ // first check input consistency
+ if (inputMatToConvert.empty())
+ throw cv::Exception(-1, "RetinaImpl cannot be applied, input buffer is empty", "RetinaImpl::run", "RetinaImpl.h", 0);
+
+ // retreive color mode from image input
+ int imageNumberOfChannels = inputMatToConvert.channels();
+
+ // convert to float AND fill the valarray buffer
+ typedef float T; // define here the target pixel format, here, float
+ const int dsttype = DataType<T>::depth; // output buffer is float format
+
+ const unsigned int nbPixels=inputMat.getMat().rows*inputMat.getMat().cols;
+ const unsigned int doubleNBpixels=inputMat.getMat().rows*inputMat.getMat().cols*2;
+
+ if(imageNumberOfChannels==4)
+ {
+ // create a cv::Mat table (for RGBA planes)
+ cv::Mat planes[4] =
+ {
+ cv::Mat(inputMatToConvert.size(), dsttype, &outputValarrayMatrix[doubleNBpixels]),
+ cv::Mat(inputMatToConvert.size(), dsttype, &outputValarrayMatrix[nbPixels]),
+ cv::Mat(inputMatToConvert.size(), dsttype, &outputValarrayMatrix[0])
+ };
+ planes[3] = cv::Mat(inputMatToConvert.size(), dsttype); // last channel (alpha) does not point on the valarray (not usefull in our case)
+ // split color cv::Mat in 4 planes... it fills valarray directely
+ cv::split(Mat_<Vec<T, 4> >(inputMatToConvert), planes);
+ }
+ else if (imageNumberOfChannels==3)
+ {
+ // create a cv::Mat table (for RGB planes)
+ cv::Mat planes[] =
+ {
+ cv::Mat(inputMatToConvert.size(), dsttype, &outputValarrayMatrix[doubleNBpixels]),
+ cv::Mat(inputMatToConvert.size(), dsttype, &outputValarrayMatrix[nbPixels]),
+ cv::Mat(inputMatToConvert.size(), dsttype, &outputValarrayMatrix[0])
+ };
+ // split color cv::Mat in 3 planes... it fills valarray directely
+ cv::split(cv::Mat_<Vec<T, 3> >(inputMatToConvert), planes);
+ }
+ else if(imageNumberOfChannels==1)
+ {
+ // create a cv::Mat header for the valarray
+ cv::Mat dst(inputMatToConvert.size(), dsttype, &outputValarrayMatrix[0]);
+ inputMatToConvert.convertTo(dst, dsttype);
+ }
+ else
+ CV_Error(Error::StsUnsupportedFormat, "input image must be single channel (gray levels), bgr format (color) or bgra (color with transparency which won't be considered");
+
+ return imageNumberOfChannels>1; // return bool : false for gray level image processing, true for color mode
+}
+
+void RetinaImpl::clearBuffers() {_retinaFilter->clearAllBuffers();}
+
+void RetinaImpl::activateMovingContoursProcessing(const bool activate){_retinaFilter->activateMovingContoursProcessing(activate);}
+
+void RetinaImpl::activateContoursProcessing(const bool activate){_retinaFilter->activateContoursProcessing(activate);}
+
+}// end of namespace bioinspired
+}// end of namespace cv
--- /dev/null
- GPU_SANITY_CHECK(u, 1e-2);
- GPU_SANITY_CHECK(v, 1e-2);
+/*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) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+// * Redistribution's of source code must retain the above copyright notice,
+// this list of conditions and the following disclaimer.
+//
+// * Redistribution's in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// * The name of the copyright holders may not be used to endorse or promote products
+// derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "perf_precomp.hpp"
+#include "opencv2/legacy.hpp"
+
+using namespace std;
+using namespace testing;
+using namespace perf;
+
+//////////////////////////////////////////////////////
+// InterpolateFrames
+
+typedef pair<string, string> pair_string;
+
+DEF_PARAM_TEST_1(ImagePair, pair_string);
+
+PERF_TEST_P(ImagePair, InterpolateFrames,
+ Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
+{
+ cv::Mat frame0 = readImage(GetParam().first, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame0.empty());
+
+ cv::Mat frame1 = readImage(GetParam().second, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame1.empty());
+
+ frame0.convertTo(frame0, CV_32FC1, 1.0 / 255.0);
+ frame1.convertTo(frame1, CV_32FC1, 1.0 / 255.0);
+
+ if (PERF_RUN_GPU())
+ {
+ const cv::gpu::GpuMat d_frame0(frame0);
+ const cv::gpu::GpuMat d_frame1(frame1);
+ cv::gpu::GpuMat d_fu, d_fv;
+ cv::gpu::GpuMat d_bu, d_bv;
+
+ cv::gpu::BroxOpticalFlow d_flow(0.197f /*alpha*/, 50.0f /*gamma*/, 0.8f /*scale_factor*/,
+ 10 /*inner_iterations*/, 77 /*outer_iterations*/, 10 /*solver_iterations*/);
+
+ d_flow(d_frame0, d_frame1, d_fu, d_fv);
+ d_flow(d_frame1, d_frame0, d_bu, d_bv);
+
+ cv::gpu::GpuMat newFrame;
+ cv::gpu::GpuMat d_buf;
+
+ TEST_CYCLE() cv::gpu::interpolateFrames(d_frame0, d_frame1, d_fu, d_fv, d_bu, d_bv, 0.5f, newFrame, d_buf);
+
+ GPU_SANITY_CHECK(newFrame, 1e-4);
+ }
+ else
+ {
+ FAIL_NO_CPU();
+ }
+}
+
+//////////////////////////////////////////////////////
+// CreateOpticalFlowNeedleMap
+
+PERF_TEST_P(ImagePair, CreateOpticalFlowNeedleMap,
+ Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
+{
+ cv::Mat frame0 = readImage(GetParam().first, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame0.empty());
+
+ cv::Mat frame1 = readImage(GetParam().second, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame1.empty());
+
+ frame0.convertTo(frame0, CV_32FC1, 1.0 / 255.0);
+ frame1.convertTo(frame1, CV_32FC1, 1.0 / 255.0);
+
+ if (PERF_RUN_GPU())
+ {
+ const cv::gpu::GpuMat d_frame0(frame0);
+ const cv::gpu::GpuMat d_frame1(frame1);
+ cv::gpu::GpuMat u;
+ cv::gpu::GpuMat v;
+
+ cv::gpu::BroxOpticalFlow d_flow(0.197f /*alpha*/, 50.0f /*gamma*/, 0.8f /*scale_factor*/,
+ 10 /*inner_iterations*/, 77 /*outer_iterations*/, 10 /*solver_iterations*/);
+
+ d_flow(d_frame0, d_frame1, u, v);
+
+ cv::gpu::GpuMat vertex, colors;
+
+ TEST_CYCLE() cv::gpu::createOpticalFlowNeedleMap(u, v, vertex, colors);
+
+ GPU_SANITY_CHECK(vertex, 1e-6);
+ GPU_SANITY_CHECK(colors);
+ }
+ else
+ {
+ FAIL_NO_CPU();
+ }
+}
+
+//////////////////////////////////////////////////////
+// BroxOpticalFlow
+
+PERF_TEST_P(ImagePair, BroxOpticalFlow,
+ Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
+{
+ declare.time(300);
+
+ cv::Mat frame0 = readImage(GetParam().first, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame0.empty());
+
+ cv::Mat frame1 = readImage(GetParam().second, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame1.empty());
+
+ frame0.convertTo(frame0, CV_32FC1, 1.0 / 255.0);
+ frame1.convertTo(frame1, CV_32FC1, 1.0 / 255.0);
+
+ if (PERF_RUN_GPU())
+ {
+ const cv::gpu::GpuMat d_frame0(frame0);
+ const cv::gpu::GpuMat d_frame1(frame1);
+ cv::gpu::GpuMat u;
+ cv::gpu::GpuMat v;
+
+ cv::gpu::BroxOpticalFlow d_flow(0.197f /*alpha*/, 50.0f /*gamma*/, 0.8f /*scale_factor*/,
+ 10 /*inner_iterations*/, 77 /*outer_iterations*/, 10 /*solver_iterations*/);
+
+ TEST_CYCLE() d_flow(d_frame0, d_frame1, u, v);
+
+ GPU_SANITY_CHECK(u, 1e-1);
+ GPU_SANITY_CHECK(v, 1e-1);
+ }
+ else
+ {
+ FAIL_NO_CPU();
+ }
+}
+
+//////////////////////////////////////////////////////
+// PyrLKOpticalFlowSparse
+
+DEF_PARAM_TEST(ImagePair_Gray_NPts_WinSz_Levels_Iters, pair_string, bool, int, int, int, int);
+
+PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, PyrLKOpticalFlowSparse,
+ Combine(Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")),
+ Bool(),
+ Values(8000),
+ Values(21),
+ Values(1, 3),
+ Values(1, 30)))
+{
+ declare.time(20.0);
+
+ const pair_string imagePair = GET_PARAM(0);
+ const bool useGray = GET_PARAM(1);
+ const int points = GET_PARAM(2);
+ const int winSize = GET_PARAM(3);
+ const int levels = GET_PARAM(4);
+ const int iters = GET_PARAM(5);
+
+ const cv::Mat frame0 = readImage(imagePair.first, useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
+ ASSERT_FALSE(frame0.empty());
+
+ const cv::Mat frame1 = readImage(imagePair.second, useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
+ ASSERT_FALSE(frame1.empty());
+
+ cv::Mat gray_frame;
+ if (useGray)
+ gray_frame = frame0;
+ else
+ cv::cvtColor(frame0, gray_frame, cv::COLOR_BGR2GRAY);
+
+ cv::Mat pts;
+ cv::goodFeaturesToTrack(gray_frame, pts, points, 0.01, 0.0);
+
+ if (PERF_RUN_GPU())
+ {
+ const cv::gpu::GpuMat d_pts(pts.reshape(2, 1));
+
+ cv::gpu::PyrLKOpticalFlow d_pyrLK;
+ d_pyrLK.winSize = cv::Size(winSize, winSize);
+ d_pyrLK.maxLevel = levels - 1;
+ d_pyrLK.iters = iters;
+
+ const cv::gpu::GpuMat d_frame0(frame0);
+ const cv::gpu::GpuMat d_frame1(frame1);
+ cv::gpu::GpuMat nextPts;
+ cv::gpu::GpuMat status;
+
+ TEST_CYCLE() d_pyrLK.sparse(d_frame0, d_frame1, d_pts, nextPts, status);
+
+ GPU_SANITY_CHECK(nextPts);
+ GPU_SANITY_CHECK(status);
+ }
+ else
+ {
+ cv::Mat nextPts;
+ cv::Mat status;
+
+ TEST_CYCLE()
+ {
+ cv::calcOpticalFlowPyrLK(frame0, frame1, pts, nextPts, status, cv::noArray(),
+ cv::Size(winSize, winSize), levels - 1,
+ cv::TermCriteria(cv::TermCriteria::COUNT + cv::TermCriteria::EPS, iters, 0.01));
+ }
+
+ CPU_SANITY_CHECK(nextPts);
+ CPU_SANITY_CHECK(status);
+ }
+}
+
+//////////////////////////////////////////////////////
+// PyrLKOpticalFlowDense
+
+DEF_PARAM_TEST(ImagePair_WinSz_Levels_Iters, pair_string, int, int, int);
+
+PERF_TEST_P(ImagePair_WinSz_Levels_Iters, PyrLKOpticalFlowDense,
+ Combine(Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")),
+ Values(3, 5, 7, 9, 13, 17, 21),
+ Values(1, 3),
+ Values(1, 10)))
+{
+ declare.time(30);
+
+ const pair_string imagePair = GET_PARAM(0);
+ const int winSize = GET_PARAM(1);
+ const int levels = GET_PARAM(2);
+ const int iters = GET_PARAM(3);
+
+ const cv::Mat frame0 = readImage(imagePair.first, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame0.empty());
+
+ const cv::Mat frame1 = readImage(imagePair.second, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame1.empty());
+
+ if (PERF_RUN_GPU())
+ {
+ const cv::gpu::GpuMat d_frame0(frame0);
+ const cv::gpu::GpuMat d_frame1(frame1);
+ cv::gpu::GpuMat u;
+ cv::gpu::GpuMat v;
+
+ cv::gpu::PyrLKOpticalFlow d_pyrLK;
+ d_pyrLK.winSize = cv::Size(winSize, winSize);
+ d_pyrLK.maxLevel = levels - 1;
+ d_pyrLK.iters = iters;
+
+ TEST_CYCLE() d_pyrLK.dense(d_frame0, d_frame1, u, v);
+
+ GPU_SANITY_CHECK(u);
+ GPU_SANITY_CHECK(v);
+ }
+ else
+ {
+ FAIL_NO_CPU();
+ }
+}
+
+//////////////////////////////////////////////////////
+// FarnebackOpticalFlow
+
+PERF_TEST_P(ImagePair, FarnebackOpticalFlow,
+ Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
+{
+ declare.time(10);
+
+ const cv::Mat frame0 = readImage(GetParam().first, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame0.empty());
+
+ const cv::Mat frame1 = readImage(GetParam().second, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame1.empty());
+
+ const int numLevels = 5;
+ const double pyrScale = 0.5;
+ const int winSize = 13;
+ const int numIters = 10;
+ const int polyN = 5;
+ const double polySigma = 1.1;
+ const int flags = 0;
+
+ if (PERF_RUN_GPU())
+ {
+ const cv::gpu::GpuMat d_frame0(frame0);
+ const cv::gpu::GpuMat d_frame1(frame1);
+ cv::gpu::GpuMat u;
+ cv::gpu::GpuMat v;
+
+ cv::gpu::FarnebackOpticalFlow d_farneback;
+ d_farneback.numLevels = numLevels;
+ d_farneback.pyrScale = pyrScale;
+ d_farneback.winSize = winSize;
+ d_farneback.numIters = numIters;
+ d_farneback.polyN = polyN;
+ d_farneback.polySigma = polySigma;
+ d_farneback.flags = flags;
+
+ TEST_CYCLE() d_farneback(d_frame0, d_frame1, u, v);
+
+ GPU_SANITY_CHECK(u, 1e-4);
+ GPU_SANITY_CHECK(v, 1e-4);
+ }
+ else
+ {
+ cv::Mat flow;
+
+ TEST_CYCLE() cv::calcOpticalFlowFarneback(frame0, frame1, flow, pyrScale, numLevels, winSize, numIters, polyN, polySigma, flags);
+
+ CPU_SANITY_CHECK(flow);
+ }
+}
+
+//////////////////////////////////////////////////////
+// OpticalFlowDual_TVL1
+
+PERF_TEST_P(ImagePair, OpticalFlowDual_TVL1,
+ Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
+{
+ declare.time(20);
+
+ const cv::Mat frame0 = readImage(GetParam().first, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame0.empty());
+
+ const cv::Mat frame1 = readImage(GetParam().second, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame1.empty());
+
+ if (PERF_RUN_GPU())
+ {
+ const cv::gpu::GpuMat d_frame0(frame0);
+ const cv::gpu::GpuMat d_frame1(frame1);
+ cv::gpu::GpuMat u;
+ cv::gpu::GpuMat v;
+
+ cv::gpu::OpticalFlowDual_TVL1_GPU d_alg;
+
+ TEST_CYCLE() d_alg(d_frame0, d_frame1, u, v);
+
++ GPU_SANITY_CHECK(u, 1e-1);
++ GPU_SANITY_CHECK(v, 1e-1);
+ }
+ else
+ {
+ cv::Mat flow;
+
+ cv::Ptr<cv::DenseOpticalFlow> alg = cv::createOptFlow_DualTVL1();
+ alg->set("medianFiltering", 1);
+ alg->set("innerIterations", 1);
+ alg->set("outerIterations", 300);
+
+ TEST_CYCLE() alg->calc(frame0, frame1, flow);
+
+ CPU_SANITY_CHECK(flow);
+ }
+}
+
+//////////////////////////////////////////////////////
+// OpticalFlowBM
+
+void calcOpticalFlowBM(const cv::Mat& prev, const cv::Mat& curr,
+ cv::Size bSize, cv::Size shiftSize, cv::Size maxRange, int usePrevious,
+ cv::Mat& velx, cv::Mat& vely)
+{
+ cv::Size sz((curr.cols - bSize.width + shiftSize.width)/shiftSize.width, (curr.rows - bSize.height + shiftSize.height)/shiftSize.height);
+
+ velx.create(sz, CV_32FC1);
+ vely.create(sz, CV_32FC1);
+
+ CvMat cvprev = prev;
+ CvMat cvcurr = curr;
+
+ CvMat cvvelx = velx;
+ CvMat cvvely = vely;
+
+ cvCalcOpticalFlowBM(&cvprev, &cvcurr, bSize, shiftSize, maxRange, usePrevious, &cvvelx, &cvvely);
+}
+
+PERF_TEST_P(ImagePair, OpticalFlowBM,
+ Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
+{
+ declare.time(400);
+
+ const cv::Mat frame0 = readImage(GetParam().first, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame0.empty());
+
+ const cv::Mat frame1 = readImage(GetParam().second, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame1.empty());
+
+ const cv::Size block_size(16, 16);
+ const cv::Size shift_size(1, 1);
+ const cv::Size max_range(16, 16);
+
+ if (PERF_RUN_GPU())
+ {
+ const cv::gpu::GpuMat d_frame0(frame0);
+ const cv::gpu::GpuMat d_frame1(frame1);
+ cv::gpu::GpuMat u, v, buf;
+
+ TEST_CYCLE() cv::gpu::calcOpticalFlowBM(d_frame0, d_frame1, block_size, shift_size, max_range, false, u, v, buf);
+
+ GPU_SANITY_CHECK(u);
+ GPU_SANITY_CHECK(v);
+ }
+ else
+ {
+ cv::Mat u, v;
+
+ TEST_CYCLE() calcOpticalFlowBM(frame0, frame1, block_size, shift_size, max_range, false, u, v);
+
+ CPU_SANITY_CHECK(u);
+ CPU_SANITY_CHECK(v);
+ }
+}
+
+PERF_TEST_P(ImagePair, FastOpticalFlowBM,
+ Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
+{
+ declare.time(400);
+
+ const cv::Mat frame0 = readImage(GetParam().first, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame0.empty());
+
+ const cv::Mat frame1 = readImage(GetParam().second, cv::IMREAD_GRAYSCALE);
+ ASSERT_FALSE(frame1.empty());
+
+ const cv::Size block_size(16, 16);
+ const cv::Size shift_size(1, 1);
+ const cv::Size max_range(16, 16);
+
+ if (PERF_RUN_GPU())
+ {
+ const cv::gpu::GpuMat d_frame0(frame0);
+ const cv::gpu::GpuMat d_frame1(frame1);
+ cv::gpu::GpuMat u, v;
+
+ cv::gpu::FastOpticalFlowBM fastBM;
+
+ TEST_CYCLE() fastBM(d_frame0, d_frame1, u, v, max_range.width, block_size.width);
+
+ GPU_SANITY_CHECK(u, 2);
+ GPU_SANITY_CHECK(v, 2);
+ }
+ else
+ {
+ FAIL_NO_CPU();
+ }
+}
--- /dev/null
- const float l_t, const float theta)
+/*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) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+// * Redistribution's of source code must retain the above copyright notice,
+// this list of conditions and the following disclaimer.
+//
+// * Redistribution's in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// * The name of the copyright holders may not be used to endorse or promote products
+// derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#if !defined CUDA_DISABLER
+
+#include "opencv2/core/cuda/common.hpp"
+#include "opencv2/core/cuda/border_interpolate.hpp"
+#include "opencv2/core/cuda/limits.hpp"
+
+using namespace cv::gpu;
+using namespace cv::gpu::cudev;
+
+////////////////////////////////////////////////////////////
+// centeredGradient
+
+namespace tvl1flow
+{
+ __global__ void centeredGradientKernel(const PtrStepSzf src, PtrStepf dx, PtrStepf dy)
+ {
+ const int x = blockIdx.x * blockDim.x + threadIdx.x;
+ const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if (x >= src.cols || y >= src.rows)
+ return;
+
+ dx(y, x) = 0.5f * (src(y, ::min(x + 1, src.cols - 1)) - src(y, ::max(x - 1, 0)));
+ dy(y, x) = 0.5f * (src(::min(y + 1, src.rows - 1), x) - src(::max(y - 1, 0), x));
+ }
+
+ void centeredGradient(PtrStepSzf src, PtrStepSzf dx, PtrStepSzf dy)
+ {
+ const dim3 block(32, 8);
+ const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
+
+ centeredGradientKernel<<<grid, block>>>(src, dx, dy);
+ cudaSafeCall( cudaGetLastError() );
+
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
+}
+
+////////////////////////////////////////////////////////////
+// warpBackward
+
+namespace tvl1flow
+{
+ static __device__ __forceinline__ float bicubicCoeff(float x_)
+ {
+ float x = fabsf(x_);
+ if (x <= 1.0f)
+ {
+ return x * x * (1.5f * x - 2.5f) + 1.0f;
+ }
+ else if (x < 2.0f)
+ {
+ return x * (x * (-0.5f * x + 2.5f) - 4.0f) + 2.0f;
+ }
+ else
+ {
+ return 0.0f;
+ }
+ }
+
+ texture<float, cudaTextureType2D, cudaReadModeElementType> tex_I1 (false, cudaFilterModePoint, cudaAddressModeClamp);
+ texture<float, cudaTextureType2D, cudaReadModeElementType> tex_I1x(false, cudaFilterModePoint, cudaAddressModeClamp);
+ texture<float, cudaTextureType2D, cudaReadModeElementType> tex_I1y(false, cudaFilterModePoint, cudaAddressModeClamp);
+
+ __global__ void warpBackwardKernel(const PtrStepSzf I0, const PtrStepf u1, const PtrStepf u2, PtrStepf I1w, PtrStepf I1wx, PtrStepf I1wy, PtrStepf grad, PtrStepf rho)
+ {
+ const int x = blockIdx.x * blockDim.x + threadIdx.x;
+ const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if (x >= I0.cols || y >= I0.rows)
+ return;
+
+ const float u1Val = u1(y, x);
+ const float u2Val = u2(y, x);
+
+ const float wx = x + u1Val;
+ const float wy = y + u2Val;
+
+ const int xmin = ::ceilf(wx - 2.0f);
+ const int xmax = ::floorf(wx + 2.0f);
+
+ const int ymin = ::ceilf(wy - 2.0f);
+ const int ymax = ::floorf(wy + 2.0f);
+
+ float sum = 0.0f;
+ float sumx = 0.0f;
+ float sumy = 0.0f;
+ float wsum = 0.0f;
+
+ for (int cy = ymin; cy <= ymax; ++cy)
+ {
+ for (int cx = xmin; cx <= xmax; ++cx)
+ {
+ const float w = bicubicCoeff(wx - cx) * bicubicCoeff(wy - cy);
+
+ sum += w * tex2D(tex_I1 , cx, cy);
+ sumx += w * tex2D(tex_I1x, cx, cy);
+ sumy += w * tex2D(tex_I1y, cx, cy);
+
+ wsum += w;
+ }
+ }
+
+ const float coeff = 1.0f / wsum;
+
+ const float I1wVal = sum * coeff;
+ const float I1wxVal = sumx * coeff;
+ const float I1wyVal = sumy * coeff;
+
+ I1w(y, x) = I1wVal;
+ I1wx(y, x) = I1wxVal;
+ I1wy(y, x) = I1wyVal;
+
+ const float Ix2 = I1wxVal * I1wxVal;
+ const float Iy2 = I1wyVal * I1wyVal;
+
+ // store the |Grad(I1)|^2
+ grad(y, x) = Ix2 + Iy2;
+
+ // compute the constant part of the rho function
+ const float I0Val = I0(y, x);
+ rho(y, x) = I1wVal - I1wxVal * u1Val - I1wyVal * u2Val - I0Val;
+ }
+
+ void warpBackward(PtrStepSzf I0, PtrStepSzf I1, PtrStepSzf I1x, PtrStepSzf I1y, PtrStepSzf u1, PtrStepSzf u2, PtrStepSzf I1w, PtrStepSzf I1wx, PtrStepSzf I1wy, PtrStepSzf grad, PtrStepSzf rho)
+ {
+ const dim3 block(32, 8);
+ const dim3 grid(divUp(I0.cols, block.x), divUp(I0.rows, block.y));
+
+ bindTexture(&tex_I1 , I1);
+ bindTexture(&tex_I1x, I1x);
+ bindTexture(&tex_I1y, I1y);
+
+ warpBackwardKernel<<<grid, block>>>(I0, u1, u2, I1w, I1wx, I1wy, grad, rho);
+ cudaSafeCall( cudaGetLastError() );
+
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
+}
+
+////////////////////////////////////////////////////////////
+// estimateU
+
+namespace tvl1flow
+{
+ __device__ float divergence(const PtrStepf& v1, const PtrStepf& v2, int y, int x)
+ {
+ if (x > 0 && y > 0)
+ {
+ const float v1x = v1(y, x) - v1(y, x - 1);
+ const float v2y = v2(y, x) - v2(y - 1, x);
+ return v1x + v2y;
+ }
+ else
+ {
+ if (y > 0)
+ return v1(y, 0) + v2(y, 0) - v2(y - 1, 0);
+ else
+ {
+ if (x > 0)
+ return v1(0, x) - v1(0, x - 1) + v2(0, x);
+ else
+ return v1(0, 0) + v2(0, 0);
+ }
+ }
+ }
+
+ __global__ void estimateUKernel(const PtrStepSzf I1wx, const PtrStepf I1wy,
+ const PtrStepf grad, const PtrStepf rho_c,
+ const PtrStepf p11, const PtrStepf p12, const PtrStepf p21, const PtrStepf p22,
+ PtrStepf u1, PtrStepf u2, PtrStepf error,
- const float n1 = (u1OldVal - u1NewVal) * (u1OldVal - u1NewVal);
- const float n2 = (u2OldVal - u2NewVal) * (u2OldVal - u2NewVal);
- error(y, x) = n1 + n2;
++ const float l_t, const float theta, const bool calcError)
+ {
+ const int x = blockIdx.x * blockDim.x + threadIdx.x;
+ const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if (x >= I1wx.cols || y >= I1wx.rows)
+ return;
+
+ const float I1wxVal = I1wx(y, x);
+ const float I1wyVal = I1wy(y, x);
+ const float gradVal = grad(y, x);
+ const float u1OldVal = u1(y, x);
+ const float u2OldVal = u2(y, x);
+
+ const float rho = rho_c(y, x) + (I1wxVal * u1OldVal + I1wyVal * u2OldVal);
+
+ // estimate the values of the variable (v1, v2) (thresholding operator TH)
+
+ float d1 = 0.0f;
+ float d2 = 0.0f;
+
+ if (rho < -l_t * gradVal)
+ {
+ d1 = l_t * I1wxVal;
+ d2 = l_t * I1wyVal;
+ }
+ else if (rho > l_t * gradVal)
+ {
+ d1 = -l_t * I1wxVal;
+ d2 = -l_t * I1wyVal;
+ }
+ else if (gradVal > numeric_limits<float>::epsilon())
+ {
+ const float fi = -rho / gradVal;
+ d1 = fi * I1wxVal;
+ d2 = fi * I1wyVal;
+ }
+
+ const float v1 = u1OldVal + d1;
+ const float v2 = u2OldVal + d2;
+
+ // compute the divergence of the dual variable (p1, p2)
+
+ const float div_p1 = divergence(p11, p12, y, x);
+ const float div_p2 = divergence(p21, p22, y, x);
+
+ // estimate the values of the optical flow (u1, u2)
+
+ const float u1NewVal = v1 + theta * div_p1;
+ const float u2NewVal = v2 + theta * div_p2;
+
+ u1(y, x) = u1NewVal;
+ u2(y, x) = u2NewVal;
+
- float l_t, float theta)
++ if (calcError)
++ {
++ const float n1 = (u1OldVal - u1NewVal) * (u1OldVal - u1NewVal);
++ const float n2 = (u2OldVal - u2NewVal) * (u2OldVal - u2NewVal);
++ error(y, x) = n1 + n2;
++ }
+ }
+
+ void estimateU(PtrStepSzf I1wx, PtrStepSzf I1wy,
+ PtrStepSzf grad, PtrStepSzf rho_c,
+ PtrStepSzf p11, PtrStepSzf p12, PtrStepSzf p21, PtrStepSzf p22,
+ PtrStepSzf u1, PtrStepSzf u2, PtrStepSzf error,
- estimateUKernel<<<grid, block>>>(I1wx, I1wy, grad, rho_c, p11, p12, p21, p22, u1, u2, error, l_t, theta);
++ float l_t, float theta, bool calcError)
+ {
+ const dim3 block(32, 8);
+ const dim3 grid(divUp(I1wx.cols, block.x), divUp(I1wx.rows, block.y));
+
++ estimateUKernel<<<grid, block>>>(I1wx, I1wy, grad, rho_c, p11, p12, p21, p22, u1, u2, error, l_t, theta, calcError);
+ cudaSafeCall( cudaGetLastError() );
+
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
+}
+
+////////////////////////////////////////////////////////////
+// estimateDualVariables
+
+namespace tvl1flow
+{
+ __global__ void estimateDualVariablesKernel(const PtrStepSzf u1, const PtrStepf u2, PtrStepf p11, PtrStepf p12, PtrStepf p21, PtrStepf p22, const float taut)
+ {
+ const int x = blockIdx.x * blockDim.x + threadIdx.x;
+ const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if (x >= u1.cols || y >= u1.rows)
+ return;
+
+ const float u1x = u1(y, ::min(x + 1, u1.cols - 1)) - u1(y, x);
+ const float u1y = u1(::min(y + 1, u1.rows - 1), x) - u1(y, x);
+
+ const float u2x = u2(y, ::min(x + 1, u1.cols - 1)) - u2(y, x);
+ const float u2y = u2(::min(y + 1, u1.rows - 1), x) - u2(y, x);
+
+ const float g1 = ::hypotf(u1x, u1y);
+ const float g2 = ::hypotf(u2x, u2y);
+
+ const float ng1 = 1.0f + taut * g1;
+ const float ng2 = 1.0f + taut * g2;
+
+ p11(y, x) = (p11(y, x) + taut * u1x) / ng1;
+ p12(y, x) = (p12(y, x) + taut * u1y) / ng1;
+ p21(y, x) = (p21(y, x) + taut * u2x) / ng2;
+ p22(y, x) = (p22(y, x) + taut * u2y) / ng2;
+ }
+
+ void estimateDualVariables(PtrStepSzf u1, PtrStepSzf u2, PtrStepSzf p11, PtrStepSzf p12, PtrStepSzf p21, PtrStepSzf p22, float taut)
+ {
+ const dim3 block(32, 8);
+ const dim3 grid(divUp(u1.cols, block.x), divUp(u1.rows, block.y));
+
+ estimateDualVariablesKernel<<<grid, block>>>(u1, u2, p11, p12, p21, p22, taut);
+ cudaSafeCall( cudaGetLastError() );
+
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
+}
+
+#endif // !defined CUDA_DISABLER
--- /dev/null
- float l_t, float theta);
+/*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) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+// * Redistribution's of source code must retain the above copyright notice,
+// this list of conditions and the following disclaimer.
+//
+// * Redistribution's in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// * The name of the copyright holders may not be used to endorse or promote products
+// derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "precomp.hpp"
+
+#if !defined HAVE_CUDA || defined(CUDA_DISABLER)
+
+cv::gpu::OpticalFlowDual_TVL1_GPU::OpticalFlowDual_TVL1_GPU() { throw_no_cuda(); }
+void cv::gpu::OpticalFlowDual_TVL1_GPU::operator ()(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&) { throw_no_cuda(); }
+void cv::gpu::OpticalFlowDual_TVL1_GPU::collectGarbage() {}
+void cv::gpu::OpticalFlowDual_TVL1_GPU::procOneScale(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&) { throw_no_cuda(); }
+
+#else
+
+using namespace cv;
+using namespace cv::gpu;
+
+cv::gpu::OpticalFlowDual_TVL1_GPU::OpticalFlowDual_TVL1_GPU()
+{
+ tau = 0.25;
+ lambda = 0.15;
+ theta = 0.3;
+ nscales = 5;
+ warps = 5;
+ epsilon = 0.01;
+ iterations = 300;
+ scaleStep = 0.8;
+ useInitialFlow = false;
+}
+
+void cv::gpu::OpticalFlowDual_TVL1_GPU::operator ()(const GpuMat& I0, const GpuMat& I1, GpuMat& flowx, GpuMat& flowy)
+{
+ CV_Assert( I0.type() == CV_8UC1 || I0.type() == CV_32FC1 );
+ CV_Assert( I0.size() == I1.size() );
+ CV_Assert( I0.type() == I1.type() );
+ CV_Assert( !useInitialFlow || (flowx.size() == I0.size() && flowx.type() == CV_32FC1 && flowy.size() == flowx.size() && flowy.type() == flowx.type()) );
+ CV_Assert( nscales > 0 );
+
+ // allocate memory for the pyramid structure
+ I0s.resize(nscales);
+ I1s.resize(nscales);
+ u1s.resize(nscales);
+ u2s.resize(nscales);
+
+ I0.convertTo(I0s[0], CV_32F, I0.depth() == CV_8U ? 1.0 : 255.0);
+ I1.convertTo(I1s[0], CV_32F, I1.depth() == CV_8U ? 1.0 : 255.0);
+
+ if (!useInitialFlow)
+ {
+ flowx.create(I0.size(), CV_32FC1);
+ flowy.create(I0.size(), CV_32FC1);
+ }
+
+ u1s[0] = flowx;
+ u2s[0] = flowy;
+
+ I1x_buf.create(I0.size(), CV_32FC1);
+ I1y_buf.create(I0.size(), CV_32FC1);
+
+ I1w_buf.create(I0.size(), CV_32FC1);
+ I1wx_buf.create(I0.size(), CV_32FC1);
+ I1wy_buf.create(I0.size(), CV_32FC1);
+
+ grad_buf.create(I0.size(), CV_32FC1);
+ rho_c_buf.create(I0.size(), CV_32FC1);
+
+ p11_buf.create(I0.size(), CV_32FC1);
+ p12_buf.create(I0.size(), CV_32FC1);
+ p21_buf.create(I0.size(), CV_32FC1);
+ p22_buf.create(I0.size(), CV_32FC1);
+
+ diff_buf.create(I0.size(), CV_32FC1);
+
+ // create the scales
+ for (int s = 1; s < nscales; ++s)
+ {
+ gpu::resize(I0s[s-1], I0s[s], Size(), scaleStep, scaleStep);
+ gpu::resize(I1s[s-1], I1s[s], Size(), scaleStep, scaleStep);
+
+ if (I0s[s].cols < 16 || I0s[s].rows < 16)
+ {
+ nscales = s;
+ break;
+ }
+
+ if (useInitialFlow)
+ {
+ gpu::resize(u1s[s-1], u1s[s], Size(), scaleStep, scaleStep);
+ gpu::resize(u2s[s-1], u2s[s], Size(), scaleStep, scaleStep);
+
+ gpu::multiply(u1s[s], Scalar::all(scaleStep), u1s[s]);
+ gpu::multiply(u2s[s], Scalar::all(scaleStep), u2s[s]);
+ }
+ else
+ {
+ u1s[s].create(I0s[s].size(), CV_32FC1);
+ u2s[s].create(I0s[s].size(), CV_32FC1);
+ }
+ }
+
+ if (!useInitialFlow)
+ {
+ u1s[nscales-1].setTo(Scalar::all(0));
+ u2s[nscales-1].setTo(Scalar::all(0));
+ }
+
+ // pyramidal structure for computing the optical flow
+ for (int s = nscales - 1; s >= 0; --s)
+ {
+ // compute the optical flow at the current scale
+ procOneScale(I0s[s], I1s[s], u1s[s], u2s[s]);
+
+ // if this was the last scale, finish now
+ if (s == 0)
+ break;
+
+ // otherwise, upsample the optical flow
+
+ // zoom the optical flow for the next finer scale
+ gpu::resize(u1s[s], u1s[s - 1], I0s[s - 1].size());
+ gpu::resize(u2s[s], u2s[s - 1], I0s[s - 1].size());
+
+ // scale the optical flow with the appropriate zoom factor
+ gpu::multiply(u1s[s - 1], Scalar::all(1/scaleStep), u1s[s - 1]);
+ gpu::multiply(u2s[s - 1], Scalar::all(1/scaleStep), u2s[s - 1]);
+ }
+}
+
+namespace tvl1flow
+{
+ void centeredGradient(PtrStepSzf src, PtrStepSzf dx, PtrStepSzf dy);
+ void warpBackward(PtrStepSzf I0, PtrStepSzf I1, PtrStepSzf I1x, PtrStepSzf I1y, PtrStepSzf u1, PtrStepSzf u2, PtrStepSzf I1w, PtrStepSzf I1wx, PtrStepSzf I1wy, PtrStepSzf grad, PtrStepSzf rho);
+ void estimateU(PtrStepSzf I1wx, PtrStepSzf I1wy,
+ PtrStepSzf grad, PtrStepSzf rho_c,
+ PtrStepSzf p11, PtrStepSzf p12, PtrStepSzf p21, PtrStepSzf p22,
+ PtrStepSzf u1, PtrStepSzf u2, PtrStepSzf error,
- estimateU(I1wx, I1wy, grad, rho_c, p11, p12, p21, p22, u1, u2, diff, l_t, static_cast<float>(theta));
++ float l_t, float theta, bool calcError);
+ void estimateDualVariables(PtrStepSzf u1, PtrStepSzf u2, PtrStepSzf p11, PtrStepSzf p12, PtrStepSzf p21, PtrStepSzf p22, float taut);
+}
+
+void cv::gpu::OpticalFlowDual_TVL1_GPU::procOneScale(const GpuMat& I0, const GpuMat& I1, GpuMat& u1, GpuMat& u2)
+{
+ using namespace tvl1flow;
+
+ const double scaledEpsilon = epsilon * epsilon * I0.size().area();
+
+ CV_DbgAssert( I1.size() == I0.size() );
+ CV_DbgAssert( I1.type() == I0.type() );
+ CV_DbgAssert( u1.size() == I0.size() );
+ CV_DbgAssert( u2.size() == u1.size() );
+
+ GpuMat I1x = I1x_buf(Rect(0, 0, I0.cols, I0.rows));
+ GpuMat I1y = I1y_buf(Rect(0, 0, I0.cols, I0.rows));
+ centeredGradient(I1, I1x, I1y);
+
+ GpuMat I1w = I1w_buf(Rect(0, 0, I0.cols, I0.rows));
+ GpuMat I1wx = I1wx_buf(Rect(0, 0, I0.cols, I0.rows));
+ GpuMat I1wy = I1wy_buf(Rect(0, 0, I0.cols, I0.rows));
+
+ GpuMat grad = grad_buf(Rect(0, 0, I0.cols, I0.rows));
+ GpuMat rho_c = rho_c_buf(Rect(0, 0, I0.cols, I0.rows));
+
+ GpuMat p11 = p11_buf(Rect(0, 0, I0.cols, I0.rows));
+ GpuMat p12 = p12_buf(Rect(0, 0, I0.cols, I0.rows));
+ GpuMat p21 = p21_buf(Rect(0, 0, I0.cols, I0.rows));
+ GpuMat p22 = p22_buf(Rect(0, 0, I0.cols, I0.rows));
+ p11.setTo(Scalar::all(0));
+ p12.setTo(Scalar::all(0));
+ p21.setTo(Scalar::all(0));
+ p22.setTo(Scalar::all(0));
+
+ GpuMat diff = diff_buf(Rect(0, 0, I0.cols, I0.rows));
+
+ const float l_t = static_cast<float>(lambda * theta);
+ const float taut = static_cast<float>(tau / theta);
+
+ for (int warpings = 0; warpings < warps; ++warpings)
+ {
+ warpBackward(I0, I1, I1x, I1y, u1, u2, I1w, I1wx, I1wy, grad, rho_c);
+
+ double error = std::numeric_limits<double>::max();
++ double prevError = 0.0;
+ for (int n = 0; error > scaledEpsilon && n < iterations; ++n)
+ {
- if (epsilon > 0)
++ // some tweaks to make sum operation less frequently
++ bool calcError = (epsilon > 0) && (n & 0x1) && (prevError < scaledEpsilon);
+
++ estimateU(I1wx, I1wy, grad, rho_c, p11, p12, p21, p22, u1, u2, diff, l_t, static_cast<float>(theta), calcError);
++
++ if (calcError)
++ {
+ error = gpu::sum(diff, norm_buf)[0];
++ prevError = error;
++ }
++ else
++ {
++ error = std::numeric_limits<double>::max();
++ prevError -= scaledEpsilon;
++ }
+
+ estimateDualVariables(u1, u2, p11, p12, p21, p22, taut);
+ }
+ }
+}
+
+void cv::gpu::OpticalFlowDual_TVL1_GPU::collectGarbage()
+{
+ I0s.clear();
+ I1s.clear();
+ u1s.clear();
+ u2s.clear();
+
+ I1x_buf.release();
+ I1y_buf.release();
+
+ I1w_buf.release();
+ I1wx_buf.release();
+ I1wy_buf.release();
+
+ grad_buf.release();
+ rho_c_buf.release();
+
+ p11_buf.release();
+ p12_buf.release();
+ p21_buf.release();
+ p22_buf.release();
+
+ diff_buf.release();
+ norm_buf.release();
+}
+
+#endif // !defined HAVE_CUDA || defined(CUDA_DISABLER)
#include "precomp.hpp"
#include <limits>
-#include <iostream>
+
+#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n))
+ #if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
+ #define MAX_IPP8u 255
+ #define MAX_IPP16u 65535
+ #define MAX_IPP32f 1.0
+ static IppStatus sts = ippInit();
+ #endif
+
namespace cv
{
parallel_for_(Range(0, src.rows), CvtColorLoop_Invoker<Cvt>(src, dst, cvt), src.total()/(double)(1<<16) );
}
- IPPColor2GrayFunctor(ippiColor2GrayFunc _func) : func(_func)
+ #if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
+ typedef IppStatus (CV_STDCALL* ippiReorderFunc)(const void *, int, void *, int, IppiSize, const int *);
+ typedef IppStatus (CV_STDCALL* ippiGeneralFunc)(const void *, int, void *, int, IppiSize);
+ typedef IppStatus (CV_STDCALL* ippiColor2GrayFunc)(const void *, int, void *, int, IppiSize, const Ipp32f *);
+
+ template <typename Cvt>
+ class CvtColorIPPLoop_Invoker : public ParallelLoopBody
+ {
+ public:
+
+ CvtColorIPPLoop_Invoker(const Mat& _src, Mat& _dst, const Cvt& _cvt, bool *_ok) :
+ ParallelLoopBody(), src(_src), dst(_dst), cvt(_cvt), ok(_ok)
+ {
+ *ok = true;
+ }
+
+ virtual void operator()(const Range& range) const
+ {
+ const void *yS = src.ptr<uchar>(range.start);
+ void *yD = dst.ptr<uchar>(range.start);
+ if( cvt(yS, (int)src.step[0], yD, (int)dst.step[0], src.cols, range.end - range.start) < 0 )
+ *ok = false;
+ }
+
+ private:
+ const Mat& src;
+ Mat& dst;
+ const Cvt& cvt;
+ bool *ok;
+
+ const CvtColorIPPLoop_Invoker& operator= (const CvtColorIPPLoop_Invoker&);
+ };
+
+ template <typename Cvt>
+ bool CvtColorIPPLoop(const Mat& src, Mat& dst, const Cvt& cvt)
+ {
+ bool ok;
+ parallel_for_(Range(0, src.rows), CvtColorIPPLoop_Invoker<Cvt>(src, dst, cvt, &ok), src.total()/(double)(1<<16) );
+ return ok;
+ }
+
+ template <typename Cvt>
+ bool CvtColorIPPLoopCopy(Mat& src, Mat& dst, const Cvt& cvt)
+ {
+ Mat temp;
+ Mat &source = src;
+ if( src.data == dst.data )
+ {
+ src.copyTo(temp);
+ source = temp;
+ }
+ bool ok;
+ parallel_for_(Range(0, source.rows), CvtColorIPPLoop_Invoker<Cvt>(source, dst, cvt, &ok), source.total()/(double)(1<<16) );
+ return ok;
+ }
+
+ IppStatus __stdcall ippiSwapChannels_8u_C3C4Rf(const Ipp8u* pSrc, int srcStep, Ipp8u* pDst, int dstStep,
+ IppiSize roiSize, const int *dstOrder)
+ {
+ return ippiSwapChannels_8u_C3C4R(pSrc, srcStep, pDst, dstStep, roiSize, dstOrder, MAX_IPP8u);
+ }
+
+ IppStatus __stdcall ippiSwapChannels_16u_C3C4Rf(const Ipp16u* pSrc, int srcStep, Ipp16u* pDst, int dstStep,
+ IppiSize roiSize, const int *dstOrder)
+ {
+ return ippiSwapChannels_16u_C3C4R(pSrc, srcStep, pDst, dstStep, roiSize, dstOrder, MAX_IPP16u);
+ }
+
+ IppStatus __stdcall ippiSwapChannels_32f_C3C4Rf(const Ipp32f* pSrc, int srcStep, Ipp32f* pDst, int dstStep,
+ IppiSize roiSize, const int *dstOrder)
+ {
+ return ippiSwapChannels_32f_C3C4R(pSrc, srcStep, pDst, dstStep, roiSize, dstOrder, MAX_IPP32f);
+ }
+
+ static ippiReorderFunc ippiSwapChannelsC3C4RTab[] =
+ {
+ (ippiReorderFunc)ippiSwapChannels_8u_C3C4Rf, 0, (ippiReorderFunc)ippiSwapChannels_16u_C3C4Rf, 0,
+ 0, (ippiReorderFunc)ippiSwapChannels_32f_C3C4Rf, 0, 0
+ };
+
+ static ippiGeneralFunc ippiCopyAC4C3RTab[] =
+ {
+ (ippiGeneralFunc)ippiCopy_8u_AC4C3R, 0, (ippiGeneralFunc)ippiCopy_16u_AC4C3R, 0,
+ 0, (ippiGeneralFunc)ippiCopy_32f_AC4C3R, 0, 0
+ };
+
+ static ippiReorderFunc ippiSwapChannelsC4C3RTab[] =
+ {
+ (ippiReorderFunc)ippiSwapChannels_8u_C4C3R, 0, (ippiReorderFunc)ippiSwapChannels_16u_C4C3R, 0,
+ 0, (ippiReorderFunc)ippiSwapChannels_32f_C4C3R, 0, 0
+ };
+
+ static ippiReorderFunc ippiSwapChannelsC3RTab[] =
+ {
+ (ippiReorderFunc)ippiSwapChannels_8u_C3R, 0, (ippiReorderFunc)ippiSwapChannels_16u_C3R, 0,
+ 0, (ippiReorderFunc)ippiSwapChannels_32f_C3R, 0, 0
+ };
+
+ static ippiReorderFunc ippiSwapChannelsC4RTab[] =
+ {
+ (ippiReorderFunc)ippiSwapChannels_8u_AC4R, 0, (ippiReorderFunc)ippiSwapChannels_16u_AC4R, 0,
+ 0, (ippiReorderFunc)ippiSwapChannels_32f_AC4R, 0, 0
+ };
+
+ static ippiColor2GrayFunc ippiColor2GrayC3Tab[] =
+ {
+ (ippiColor2GrayFunc)ippiColorToGray_8u_C3C1R, 0, (ippiColor2GrayFunc)ippiColorToGray_16u_C3C1R, 0,
+ 0, (ippiColor2GrayFunc)ippiColorToGray_32f_C3C1R, 0, 0
+ };
+
+ static ippiColor2GrayFunc ippiColor2GrayC4Tab[] =
+ {
+ (ippiColor2GrayFunc)ippiColorToGray_8u_AC4C1R, 0, (ippiColor2GrayFunc)ippiColorToGray_16u_AC4C1R, 0,
+ 0, (ippiColor2GrayFunc)ippiColorToGray_32f_AC4C1R, 0, 0
+ };
+
+ static ippiGeneralFunc ippiRGB2GrayC3Tab[] =
+ {
+ (ippiGeneralFunc)ippiRGBToGray_8u_C3C1R, 0, (ippiGeneralFunc)ippiRGBToGray_16u_C3C1R, 0,
+ 0, (ippiGeneralFunc)ippiRGBToGray_32f_C3C1R, 0, 0
+ };
+
+ static ippiGeneralFunc ippiRGB2GrayC4Tab[] =
+ {
+ (ippiGeneralFunc)ippiRGBToGray_8u_AC4C1R, 0, (ippiGeneralFunc)ippiRGBToGray_16u_AC4C1R, 0,
+ 0, (ippiGeneralFunc)ippiRGBToGray_32f_AC4C1R, 0, 0
+ };
+
+ static ippiGeneralFunc ippiCopyP3C3RTab[] =
+ {
+ (ippiGeneralFunc)ippiCopy_8u_P3C3R, 0, (ippiGeneralFunc)ippiCopy_16u_P3C3R, 0,
+ 0, (ippiGeneralFunc)ippiCopy_32f_P3C3R, 0, 0
+ };
+
+ static ippiGeneralFunc ippiRGB2XYZTab[] =
+ {
+ (ippiGeneralFunc)ippiRGBToXYZ_8u_C3R, 0, (ippiGeneralFunc)ippiRGBToXYZ_16u_C3R, 0,
+ 0, (ippiGeneralFunc)ippiRGBToXYZ_32f_C3R, 0, 0
+ };
+
+ static ippiGeneralFunc ippiXYZ2RGBTab[] =
+ {
+ (ippiGeneralFunc)ippiXYZToRGB_8u_C3R, 0, (ippiGeneralFunc)ippiXYZToRGB_16u_C3R, 0,
+ 0, (ippiGeneralFunc)ippiXYZToRGB_32f_C3R, 0, 0
+ };
+
+ static ippiGeneralFunc ippiRGB2HSVTab[] =
+ {
+ (ippiGeneralFunc)ippiRGBToHSV_8u_C3R, 0, (ippiGeneralFunc)ippiRGBToHSV_16u_C3R, 0,
+ 0, 0, 0, 0
+ };
+
+ static ippiGeneralFunc ippiHSV2RGBTab[] =
+ {
+ (ippiGeneralFunc)ippiHSVToRGB_8u_C3R, 0, (ippiGeneralFunc)ippiHSVToRGB_16u_C3R, 0,
+ 0, 0, 0, 0
+ };
+
+ static ippiGeneralFunc ippiRGB2HLSTab[] =
+ {
+ (ippiGeneralFunc)ippiRGBToHLS_8u_C3R, 0, (ippiGeneralFunc)ippiRGBToHLS_16u_C3R, 0,
+ 0, (ippiGeneralFunc)ippiRGBToHLS_32f_C3R, 0, 0
+ };
+
+ static ippiGeneralFunc ippiHLS2RGBTab[] =
+ {
+ (ippiGeneralFunc)ippiHLSToRGB_8u_C3R, 0, (ippiGeneralFunc)ippiHLSToRGB_16u_C3R, 0,
+ 0, (ippiGeneralFunc)ippiHLSToRGB_32f_C3R, 0, 0
+ };
+
+ struct IPPGeneralFunctor
+ {
+ IPPGeneralFunctor(ippiGeneralFunc _func) : func(_func){}
+ bool operator()(const void *src, int srcStep, void *dst, int dstStep, int cols, int rows) const
+ {
+ return func(src, srcStep, dst, dstStep, ippiSize(cols, rows)) >= 0;
+ }
+ private:
+ ippiGeneralFunc func;
+ };
+
+ struct IPPReorderFunctor
+ {
+ IPPReorderFunctor(ippiReorderFunc _func, int _order0, int _order1, int _order2) : func(_func)
+ {
+ order[0] = _order0;
+ order[1] = _order1;
+ order[2] = _order2;
+ order[3] = 3;
+ }
+ bool operator()(const void *src, int srcStep, void *dst, int dstStep, int cols, int rows) const
+ {
+ return func(src, srcStep, dst, dstStep, ippiSize(cols, rows), order) >= 0;
+ }
+ private:
+ ippiReorderFunc func;
+ int order[4];
+ };
+
+ struct IPPColor2GrayFunctor
+ {
- Mat temp;
++ IPPColor2GrayFunctor(ippiColor2GrayFunc _func) : func(_func)
+ {
+ coeffs[0] = 0.114f;
+ coeffs[1] = 0.587f;
+ coeffs[2] = 0.299f;
+ }
+ bool operator()(const void *src, int srcStep, void *dst, int dstStep, int cols, int rows) const
+ {
+ return func(src, srcStep, dst, dstStep, ippiSize(cols, rows), coeffs) >= 0;
+ }
+ private:
+ ippiColor2GrayFunc func;
+ Ipp32f coeffs[3];
+ };
+
+ struct IPPGray2BGRFunctor
+ {
+ IPPGray2BGRFunctor(ippiGeneralFunc _func) : func(_func){}
+ bool operator()(const void *src, int srcStep, void *dst, int dstStep, int cols, int rows) const
+ {
+ const void* srcarray[3] = { src, src, src };
+ return func(srcarray, srcStep, dst, dstStep, ippiSize(cols, rows)) >= 0;
+ }
+ private:
+ ippiGeneralFunc func;
+ };
+
+ struct IPPGray2BGRAFunctor
+ {
+ IPPGray2BGRAFunctor(ippiGeneralFunc _func1, ippiReorderFunc _func2, int _depth) : func1(_func1), func2(_func2), depth(_depth){}
+ bool operator()(const void *src, int srcStep, void *dst, int dstStep, int cols, int rows) const
+ {
+ const void* srcarray[3] = { src, src, src };
+ Mat temp(rows, cols, CV_MAKETYPE(depth, 3));
+ if(func1(srcarray, srcStep, temp.data, (int)temp.step[0], ippiSize(cols, rows)) < 0)
+ return false;
+ int order[4] = {0, 1, 2, 3};
+ return func2(temp.data, (int)temp.step[0], dst, dstStep, ippiSize(cols, rows), order) >= 0;
+ }
+ private:
+ ippiGeneralFunc func1;
+ ippiReorderFunc func2;
+ int depth;
+ };
+
+ struct IPPReorderGeneralFunctor
+ {
+ IPPReorderGeneralFunctor(ippiReorderFunc _func1, ippiGeneralFunc _func2, int _order0, int _order1, int _order2, int _depth) : func1(_func1), func2(_func2), depth(_depth)
+ {
+ order[0] = _order0;
+ order[1] = _order1;
+ order[2] = _order2;
+ order[3] = 3;
+ }
+ bool operator()(const void *src, int srcStep, void *dst, int dstStep, int cols, int rows) const
+ {
- Mat temp;
++ Mat temp;
+ temp.create(rows, cols, CV_MAKETYPE(depth, 3));
+ if(func1(src, srcStep, temp.data, (int)temp.step[0], ippiSize(cols, rows), order) < 0)
+ return false;
+ return func2(temp.data, (int)temp.step[0], dst, dstStep, ippiSize(cols, rows)) >= 0;
+ }
+ private:
+ ippiReorderFunc func1;
+ ippiGeneralFunc func2;
+ int order[4];
+ int depth;
+ };
+
+ struct IPPGeneralReorderFunctor
+ {
+ IPPGeneralReorderFunctor(ippiGeneralFunc _func1, ippiReorderFunc _func2, int _order0, int _order1, int _order2, int _depth) : func1(_func1), func2(_func2), depth(_depth)
+ {
+ order[0] = _order0;
+ order[1] = _order1;
+ order[2] = _order2;
+ order[3] = 3;
+ }
+ bool operator()(const void *src, int srcStep, void *dst, int dstStep, int cols, int rows) const
+ {
++ Mat temp;
+ temp.create(rows, cols, CV_MAKETYPE(depth, 3));
+ if(func1(src, srcStep, temp.data, (int)temp.step[0], ippiSize(cols, rows)) < 0)
+ return false;
+ return func2(temp.data, (int)temp.step[0], dst, dstStep, ippiSize(cols, rows), order) >= 0;
+ }
+ private:
+ ippiGeneralFunc func1;
+ ippiReorderFunc func2;
+ int order[4];
+ int depth;
+ };
+ #endif
+
////////////////// Various 3/4-channel to 3/4-channel RGB transformations /////////////////
template<typename _Tp> struct RGB2RGB
_dst.create( sz, CV_MAKETYPE(depth, dcn));
dst = _dst.getMat();
-
-#if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
+
++#if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
+ if( code == CV_BGR2BGRA || code == CV_RGB2RGBA)
+ {
+ if ( CvtColorIPPLoop(src, dst, IPPReorderFunctor(ippiSwapChannelsC3C4RTab[depth], 0, 1, 2)) )
+ return;
+ }
+ else if( code == CV_BGRA2BGR )
+ {
+ if ( CvtColorIPPLoop(src, dst, IPPGeneralFunctor(ippiCopyAC4C3RTab[depth])) )
+ return;
+ }
+ else if( code == CV_BGR2RGBA )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPReorderFunctor(ippiSwapChannelsC3C4RTab[depth], 2, 1, 0)) )
+ return;
+ }
+ else if( code == CV_RGBA2BGR )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPReorderFunctor(ippiSwapChannelsC4C3RTab[depth], 2, 1, 0)) )
+ return;
+ }
+ else if( code == CV_RGB2BGR )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPReorderFunctor(ippiSwapChannelsC3RTab[depth], 2, 1, 0)) )
+ return;
+ }
+ else if( code == CV_RGBA2BGRA )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPReorderFunctor(ippiSwapChannelsC4RTab[depth], 2, 1, 0)) )
+ return;
+ }
+ #endif
+
if( depth == CV_8U )
{
#ifdef HAVE_TEGRA_OPTIMIZATION
CV_Assert( scn == 3 || scn == 4 );
_dst.create(sz, CV_MAKETYPE(depth, 1));
dst = _dst.getMat();
-
+
+ #if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
+ if( code == CV_BGR2GRAY )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPColor2GrayFunctor(ippiColor2GrayC3Tab[depth])) )
+ return;
+ }
+ else if( code == CV_RGB2GRAY )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPGeneralFunctor(ippiRGB2GrayC3Tab[depth])) )
+ return;
+ }
+ else if( code == CV_BGRA2GRAY )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPColor2GrayFunctor(ippiColor2GrayC4Tab[depth])) )
+ return;
+ }
+ else if( code == CV_RGBA2GRAY )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPGeneralFunctor(ippiRGB2GrayC4Tab[depth])) )
+ return;
+ }
+ #endif
+
bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;
if( depth == CV_8U )
CV_Assert( scn == 1 && (dcn == 3 || dcn == 4));
_dst.create(sz, CV_MAKETYPE(depth, dcn));
dst = _dst.getMat();
-
+
+ #if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
+ if( code == CV_GRAY2BGR )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPGray2BGRFunctor(ippiCopyP3C3RTab[depth])) )
+ return;
- }
++ }
+ else if( code == CV_GRAY2BGRA )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPGray2BGRAFunctor(ippiCopyP3C3RTab[depth], ippiSwapChannelsC3C4RTab[depth], depth)) )
+ return;
+ }
+ #endif
+
+
if( depth == CV_8U )
{
#ifdef HAVE_TEGRA_OPTIMIZATION
_dst.create(sz, CV_MAKETYPE(depth, 3));
dst = _dst.getMat();
-
+
+ #if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
+ if( code == CV_BGR2XYZ && scn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPReorderGeneralFunctor(ippiSwapChannelsC3RTab[depth], ippiRGB2XYZTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_BGR2XYZ && scn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPReorderGeneralFunctor(ippiSwapChannelsC4C3RTab[depth], ippiRGB2XYZTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_RGB2XYZ && scn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPGeneralFunctor(ippiRGB2XYZTab[depth])) )
+ return;
- }
++ }
+ else if( code == CV_RGB2XYZ && scn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPReorderGeneralFunctor(ippiSwapChannelsC4C3RTab[depth], ippiRGB2XYZTab[depth], 0, 1, 2, depth)) )
+ return;
+ }
+ #endif
+
if( depth == CV_8U )
CvtColorLoop(src, dst, RGB2XYZ_i<uchar>(scn, bidx, 0));
else if( depth == CV_16U )
_dst.create(sz, CV_MAKETYPE(depth, dcn));
dst = _dst.getMat();
-
+
+ #if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
+ if( code == CV_XYZ2BGR && dcn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPGeneralReorderFunctor(ippiXYZ2RGBTab[depth], ippiSwapChannelsC3RTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_XYZ2BGR && dcn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPGeneralReorderFunctor(ippiXYZ2RGBTab[depth], ippiSwapChannelsC3C4RTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ if( code == CV_XYZ2RGB && dcn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPGeneralFunctor(ippiXYZ2RGBTab[depth])) )
+ return;
+ }
+ else if( code == CV_XYZ2RGB && dcn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPGeneralReorderFunctor(ippiXYZ2RGBTab[depth], ippiSwapChannelsC3C4RTab[depth], 0, 1, 2, depth)) )
+ return;
+ }
+ #endif
+
if( depth == CV_8U )
CvtColorLoop(src, dst, XYZ2RGB_i<uchar>(dcn, bidx, 0));
else if( depth == CV_16U )
_dst.create(sz, CV_MAKETYPE(depth, 3));
dst = _dst.getMat();
-
+
+ #if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
+ if( depth == CV_8U || depth == CV_16U )
+ {
+ if( code == CV_BGR2HSV_FULL && scn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPReorderGeneralFunctor(ippiSwapChannelsC3RTab[depth], ippiRGB2HSVTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_BGR2HSV_FULL && scn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPReorderGeneralFunctor(ippiSwapChannelsC4C3RTab[depth], ippiRGB2HSVTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_RGB2HSV_FULL && scn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPGeneralFunctor(ippiRGB2HSVTab[depth])) )
+ return;
- }
++ }
+ else if( code == CV_RGB2HSV_FULL && scn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPReorderGeneralFunctor(ippiSwapChannelsC4C3RTab[depth], ippiRGB2HSVTab[depth], 0, 1, 2, depth)) )
+ return;
- }
++ }
+ else if( code == CV_BGR2HLS_FULL && scn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPReorderGeneralFunctor(ippiSwapChannelsC3RTab[depth], ippiRGB2HLSTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_BGR2HLS_FULL && scn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPReorderGeneralFunctor(ippiSwapChannelsC4C3RTab[depth], ippiRGB2HLSTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_RGB2HLS_FULL && scn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPGeneralFunctor(ippiRGB2HLSTab[depth])) )
+ return;
- }
++ }
+ else if( code == CV_RGB2HLS_FULL && scn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPReorderGeneralFunctor(ippiSwapChannelsC4C3RTab[depth], ippiRGB2HLSTab[depth], 0, 1, 2, depth)) )
+ return;
+ }
+ }
+ #endif
+
if( code == CV_BGR2HSV || code == CV_RGB2HSV ||
code == CV_BGR2HSV_FULL || code == CV_RGB2HSV_FULL )
{
_dst.create(sz, CV_MAKETYPE(depth, dcn));
dst = _dst.getMat();
-
+
+ #if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
+ if( depth == CV_8U || depth == CV_16U )
+ {
+ if( code == CV_HSV2BGR_FULL && dcn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPGeneralReorderFunctor(ippiHSV2RGBTab[depth], ippiSwapChannelsC3RTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_HSV2BGR_FULL && dcn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPGeneralReorderFunctor(ippiHSV2RGBTab[depth], ippiSwapChannelsC3C4RTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_HSV2RGB_FULL && dcn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPGeneralFunctor(ippiHSV2RGBTab[depth])) )
+ return;
- }
++ }
+ else if( code == CV_HSV2RGB_FULL && dcn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPGeneralReorderFunctor(ippiHSV2RGBTab[depth], ippiSwapChannelsC3C4RTab[depth], 0, 1, 2, depth)) )
+ return;
- }
++ }
+ else if( code == CV_HLS2BGR_FULL && dcn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPGeneralReorderFunctor(ippiHLS2RGBTab[depth], ippiSwapChannelsC3RTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_HLS2BGR_FULL && dcn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPGeneralReorderFunctor(ippiHLS2RGBTab[depth], ippiSwapChannelsC3C4RTab[depth], 2, 1, 0, depth)) )
+ return;
+ }
+ else if( code == CV_HLS2RGB_FULL && dcn == 3 )
+ {
+ if( CvtColorIPPLoopCopy(src, dst, IPPGeneralFunctor(ippiHLS2RGBTab[depth])) )
+ return;
- }
++ }
+ else if( code == CV_HLS2RGB_FULL && dcn == 4 )
+ {
+ if( CvtColorIPPLoop(src, dst, IPPGeneralReorderFunctor(ippiHLS2RGBTab[depth], ippiSwapChannelsC3C4RTab[depth], 0, 1, 2, depth)) )
+ return;
+ }
+ }
+ #endif
+
if( code == CV_HSV2BGR || code == CV_HSV2RGB ||
code == CV_HSV2BGR_FULL || code == CV_HSV2RGB_FULL )
{
}
else if (RUN_PLAIN_IMPL)
{
- StereoBM bm(0, n_disp, winSize);
+ Ptr<StereoBM> bm = createStereoBM(n_disp, winSize);
- TEST_CYCLE() bm(left_image, right_image, disp);
+ TEST_CYCLE() bm->compute(left_image, right_image, disp);
-
- SANITY_CHECK(disp);
}
else
OCL_PERF_ELSE
{
ocl::oclMat oclSrc(src), oclDst(src.size(), CV_8UC4);
- TEST_CYCLE() ocl::cvtColor(oclSrc, oclDst, COLOR_RGBA2GRAY, 4);
- OCL_TEST_CYCLE() ocl::cvtColor(oclSrc, oclDst, CV_RGBA2GRAY, 4);
++ OCL_TEST_CYCLE() ocl::cvtColor(oclSrc, oclDst, COLOR_RGBA2GRAY, 4);
oclDst.download(dst);
SANITY_CHECK(dst);
{
ocl::oclMat oclSrc(src), oclTempl(templ), oclDst(dstSize, CV_32F);
- TEST_CYCLE() cv::ocl::matchTemplate(oclSrc, oclTempl, oclDst, TM_CCORR);
- OCL_TEST_CYCLE() cv::ocl::matchTemplate(oclSrc, oclTempl, oclDst, CV_TM_CCORR);
++ OCL_TEST_CYCLE() cv::ocl::matchTemplate(oclSrc, oclTempl, oclDst, TM_CCORR);
oclDst.download(dst);
{
ocl::oclMat oclSrc(src), oclTempl(templ), oclDst(dstSize, CV_8UC1);
- TEST_CYCLE() cv::ocl::matchTemplate(oclSrc, oclTempl, oclDst, TM_CCORR_NORMED);
- OCL_TEST_CYCLE() cv::ocl::matchTemplate(oclSrc, oclTempl, oclDst, CV_TM_CCORR_NORMED);
++ OCL_TEST_CYCLE() cv::ocl::matchTemplate(oclSrc, oclTempl, oclDst, TM_CCORR_NORMED);
oclDst.download(dst);
#include <exception>
#include <stdio.h>
-#include "opencv2/imgproc/imgproc.hpp"
-#include "opencv2/imgproc/imgproc_c.h"
-#include "opencv2/core/core_c.h"
-#include "opencv2/objdetect/objdetect.hpp"
-#include "opencv2/ocl/ocl.hpp"
+#undef OPENCV_NOSTL
+
+#include "opencv2/imgproc.hpp"
+#include "opencv2/objdetect/objdetect_c.h"
+#include "opencv2/ocl.hpp"
-#include "opencv2/core/internal.hpp"
+#include "opencv2/core/utility.hpp"
+#include "opencv2/core/private.hpp"
- //#include "opencv2/highgui.hpp"
-
#define __ATI__
#if defined (HAVE_OPENCL)
#endif
#include <iostream>
- #include "cvconfig.h"
- #include "opencv2/core/core.hpp"
- #include "opencv2/highgui/highgui.hpp"
- #include "opencv2/gpu/gpu.hpp"
-
- #ifdef HAVE_TBB
- # include "tbb/tbb_stddef.h"
- # if TBB_VERSION_MAJOR*100 + TBB_VERSION_MINOR >= 202
- # include "tbb/tbb.h"
- # include "tbb/task.h"
- # undef min
- # undef max
- # else
- # undef HAVE_TBB
- # endif
- #endif
+ #include <iomanip>
+
-#include "opencv2/core/core.hpp"
-#include "opencv2/highgui/highgui.hpp"
-#include "opencv2/imgproc/imgproc.hpp"
-#include "opencv2/contrib/contrib.hpp"
-#include "opencv2/gpu/gpu.hpp"
++#include "opencv2/core.hpp"
++#include "opencv2/highgui.hpp"
++#include "opencv2/imgproc.hpp"
++#include "opencv2/contrib.hpp"
++#include "opencv2/gpustereo.hpp"
+
+ using namespace std;
+ using namespace cv;
+ using namespace cv::gpu;
- #if !defined(HAVE_CUDA) || !defined(HAVE_TBB)
+ ///////////////////////////////////////////////////////////
+ // Thread
+ // OS-specific wrappers for multi-threading
- int main()
+ #ifdef WIN32
+ class Thread
{
- #if !defined(HAVE_CUDA)
- std::cout << "CUDA support is required (CMake key 'WITH_CUDA' must be true).\n";
- #endif
+ struct UserData
+ {
+ void (*func)(void* userData);
+ void* param;
+ };
+
+ static DWORD WINAPI WinThreadFunction(LPVOID lpParam)
+ {
+ UserData* userData = static_cast<UserData*>(lpParam);
+
+ userData->func(userData->param);
+
+ return 0;
+ }
+
+ UserData userData_;
+ HANDLE thread_;
+ DWORD threadId_;
+
+ public:
+ Thread(void (*func)(void* userData), void* userData)
+ {
+ userData_.func = func;
+ userData_.param = userData;
+
+ thread_ = CreateThread(
+ NULL, // default security attributes
+ 0, // use default stack size
+ WinThreadFunction, // thread function name
+ &userData_, // argument to thread function
+ 0, // use default creation flags
+ &threadId_); // returns the thread identifier
+ }
+
+ ~Thread()
+ {
+ CloseHandle(thread_);
+ }
+
+ void wait()
+ {
+ WaitForSingleObject(thread_, INFINITE);
+ }
+ };
+ #else
+ class Thread
+ {
+ struct UserData
+ {
+ void (*func)(void* userData);
+ void* param;
+ };
+
+ static void* PThreadFunction(void* lpParam)
+ {
+ UserData* userData = static_cast<UserData*>(lpParam);
+
+ userData->func(userData->param);
+
+ return 0;
+ }
+
+ pthread_t thread_;
+ UserData userData_;
+
+ public:
+ Thread(void (*func)(void* userData), void* userData)
+ {
+ userData_.func = func;
+ userData_.param = userData;
+
+ pthread_create(&thread_, NULL, PThreadFunction, &userData_);
+ }
- #if !defined(HAVE_TBB)
- std::cout << "TBB support is required (CMake key 'WITH_TBB' must be true).\n";
+ ~Thread()
+ {
+ pthread_detach(thread_);
+ }
+
+ void wait()
+ {
+ pthread_join(thread_, NULL);
+ }
+ };
#endif
- return 0;
+ ///////////////////////////////////////////////////////////
+ // StereoSingleGpu
+ // Run Stereo algorithm on single GPU
+
+ class StereoSingleGpu
+ {
+ public:
+ explicit StereoSingleGpu(int deviceId = 0);
+ ~StereoSingleGpu();
+
+ void compute(const Mat& leftFrame, const Mat& rightFrame, Mat& disparity);
+
+ private:
+ int deviceId_;
+ GpuMat d_leftFrame;
+ GpuMat d_rightFrame;
+ GpuMat d_disparity;
- Ptr<StereoBM_GPU> d_alg;
++ Ptr<gpu::StereoBM> d_alg;
+ };
+
+ StereoSingleGpu::StereoSingleGpu(int deviceId) : deviceId_(deviceId)
+ {
+ gpu::setDevice(deviceId_);
- d_alg = new StereoBM_GPU(StereoBM_GPU::BASIC_PRESET, 256);
++ d_alg = gpu::createStereoBM(256);
}
- #else
+ StereoSingleGpu::~StereoSingleGpu()
+ {
+ gpu::setDevice(deviceId_);
+ d_leftFrame.release();
+ d_rightFrame.release();
+ d_disparity.release();
+ d_alg.release();
+ }
- using namespace std;
- using namespace cv;
- using namespace cv::gpu;
+ void StereoSingleGpu::compute(const Mat& leftFrame, const Mat& rightFrame, Mat& disparity)
+ {
+ gpu::setDevice(deviceId_);
+ d_leftFrame.upload(leftFrame);
+ d_rightFrame.upload(rightFrame);
- (*d_alg)(d_leftFrame, d_rightFrame, d_disparity);
++ d_alg->compute(d_leftFrame, d_rightFrame, d_disparity);
+ d_disparity.download(disparity);
+ }
+
+ ///////////////////////////////////////////////////////////
+ // StereoMultiGpuThread
+ // Run Stereo algorithm on two GPUs using different host threads
+
+ class StereoMultiGpuThread
+ {
+ public:
+ StereoMultiGpuThread();
+ ~StereoMultiGpuThread();
+
+ void compute(const Mat& leftFrame, const Mat& rightFrame, Mat& disparity);
+
+ private:
+ GpuMat d_leftFrames[2];
+ GpuMat d_rightFrames[2];
+ GpuMat d_disparities[2];
- Ptr<StereoBM_GPU> d_algs[2];
++ Ptr<gpu::StereoBM> d_algs[2];
+
+ struct StereoLaunchData
+ {
+ int deviceId;
+ Mat leftFrame;
+ Mat rightFrame;
+ Mat disparity;
+ GpuMat* d_leftFrame;
+ GpuMat* d_rightFrame;
+ GpuMat* d_disparity;
- Ptr<StereoBM_GPU> d_alg;
++ Ptr<gpu::StereoBM> d_alg;
+ };
+
+ static void launchGpuStereoAlg(void* userData);
+ };
+
+ StereoMultiGpuThread::StereoMultiGpuThread()
+ {
+ gpu::setDevice(0);
- d_algs[0] = new StereoBM_GPU(StereoBM_GPU::BASIC_PRESET, 256);
++ d_algs[0] = gpu::createStereoBM(256);
+
+ gpu::setDevice(1);
- d_algs[1] = new StereoBM_GPU(StereoBM_GPU::BASIC_PRESET, 256);
++ d_algs[1] = gpu::createStereoBM(256);
+ }
+
+ StereoMultiGpuThread::~StereoMultiGpuThread()
+ {
+ gpu::setDevice(0);
+ d_leftFrames[0].release();
+ d_rightFrames[0].release();
+ d_disparities[0].release();
+ d_algs[0].release();
+
+ gpu::setDevice(1);
+ d_leftFrames[1].release();
+ d_rightFrames[1].release();
+ d_disparities[1].release();
+ d_algs[1].release();
+ }
- struct Worker { void operator()(int device_id) const; };
+ void StereoMultiGpuThread::compute(const Mat& leftFrame, const Mat& rightFrame, Mat& disparity)
+ {
+ disparity.create(leftFrame.size(), CV_8UC1);
+
+ // Split input data onto two parts for each GPUs.
+ // We add small border for each part,
+ // because original algorithm doesn't calculate disparity on image borders.
+ // With such padding we will get output in the middle of final result.
+
+ StereoLaunchData launchDatas[2];
+
+ launchDatas[0].deviceId = 0;
+ launchDatas[0].leftFrame = leftFrame.rowRange(0, leftFrame.rows / 2 + 32);
+ launchDatas[0].rightFrame = rightFrame.rowRange(0, rightFrame.rows / 2 + 32);
+ launchDatas[0].disparity = disparity.rowRange(0, leftFrame.rows / 2);
+ launchDatas[0].d_leftFrame = &d_leftFrames[0];
+ launchDatas[0].d_rightFrame = &d_rightFrames[0];
+ launchDatas[0].d_disparity = &d_disparities[0];
+ launchDatas[0].d_alg = d_algs[0];
+
+ launchDatas[1].deviceId = 1;
+ launchDatas[1].leftFrame = leftFrame.rowRange(leftFrame.rows / 2 - 32, leftFrame.rows);
+ launchDatas[1].rightFrame = rightFrame.rowRange(leftFrame.rows / 2 - 32, leftFrame.rows);
+ launchDatas[1].disparity = disparity.rowRange(leftFrame.rows / 2, leftFrame.rows);
+ launchDatas[1].d_leftFrame = &d_leftFrames[1];
+ launchDatas[1].d_rightFrame = &d_rightFrames[1];
+ launchDatas[1].d_disparity = &d_disparities[1];
+ launchDatas[1].d_alg = d_algs[1];
- // GPUs data
- GpuMat d_left[2];
- GpuMat d_right[2];
- Ptr<gpu::StereoBM> bm[2];
- GpuMat d_result[2];
+ Thread thread0(launchGpuStereoAlg, &launchDatas[0]);
+ Thread thread1(launchGpuStereoAlg, &launchDatas[1]);
- static void printHelp()
+ thread0.wait();
+ thread1.wait();
+ }
+
+ void StereoMultiGpuThread::launchGpuStereoAlg(void* userData)
{
- std::cout << "Usage: stereo_multi_gpu --left <image> --right <image>\n";
+ StereoLaunchData* data = static_cast<StereoLaunchData*>(userData);
+
+ gpu::setDevice(data->deviceId);
+ data->d_leftFrame->upload(data->leftFrame);
+ data->d_rightFrame->upload(data->rightFrame);
- (*data->d_alg)(*data->d_leftFrame, *data->d_rightFrame, *data->d_disparity);
++ data->d_alg->compute(*data->d_leftFrame, *data->d_rightFrame, *data->d_disparity);
+
+ if (data->deviceId == 0)
+ data->d_disparity->rowRange(0, data->d_disparity->rows - 32).download(data->disparity);
+ else
+ data->d_disparity->rowRange(32, data->d_disparity->rows).download(data->disparity);
}
- Ptr<StereoBM_GPU> d_algs[2];
+ ///////////////////////////////////////////////////////////
+ // StereoMultiGpuStream
+ // Run Stereo algorithm on two GPUs from single host thread using async API
+
+ class StereoMultiGpuStream
+ {
+ public:
+ StereoMultiGpuStream();
+ ~StereoMultiGpuStream();
+
+ void compute(const CudaMem& leftFrame, const CudaMem& rightFrame, CudaMem& disparity);
+
+ private:
+ GpuMat d_leftFrames[2];
+ GpuMat d_rightFrames[2];
+ GpuMat d_disparities[2];
- d_algs[0] = new StereoBM_GPU(StereoBM_GPU::BASIC_PRESET, 256);
++ Ptr<gpu::StereoBM> d_algs[2];
+ Ptr<Stream> streams[2];
+ };
+
+ StereoMultiGpuStream::StereoMultiGpuStream()
+ {
+ gpu::setDevice(0);
- d_algs[1] = new StereoBM_GPU(StereoBM_GPU::BASIC_PRESET, 256);
++ d_algs[0] = gpu::createStereoBM(256);
+ streams[0] = new Stream;
+
+ gpu::setDevice(1);
- streams[0]->enqueueUpload(leftFrameHdr.rowRange(0, leftFrame.rows / 2 + 32), d_leftFrames[0]);
- streams[0]->enqueueUpload(rightFrameHdr.rowRange(0, leftFrame.rows / 2 + 32), d_rightFrames[0]);
- (*d_algs[0])(d_leftFrames[0], d_rightFrames[0], d_disparities[0], *streams[0]);
- streams[0]->enqueueDownload(d_disparities[0].rowRange(0, leftFrame.rows / 2), disparityPart0);
++ d_algs[1] = gpu::createStereoBM(256);
+ streams[1] = new Stream;
+ }
+
+ StereoMultiGpuStream::~StereoMultiGpuStream()
+ {
+ gpu::setDevice(0);
+ d_leftFrames[0].release();
+ d_rightFrames[0].release();
+ d_disparities[0].release();
+ d_algs[0].release();
+ streams[0].release();
+
+ gpu::setDevice(1);
+ d_leftFrames[1].release();
+ d_rightFrames[1].release();
+ d_disparities[1].release();
+ d_algs[1].release();
+ streams[1].release();
+ }
+
+ void StereoMultiGpuStream::compute(const CudaMem& leftFrame, const CudaMem& rightFrame, CudaMem& disparity)
+ {
+ disparity.create(leftFrame.size(), CV_8UC1);
+
+ // Split input data onto two parts for each GPUs.
+ // We add small border for each part,
+ // because original algorithm doesn't calculate disparity on image borders.
+ // With such padding we will get output in the middle of final result.
+
+ Mat leftFrameHdr = leftFrame.createMatHeader();
+ Mat rightFrameHdr = rightFrame.createMatHeader();
+ Mat disparityHdr = disparity.createMatHeader();
+ Mat disparityPart0 = disparityHdr.rowRange(0, leftFrame.rows / 2);
+ Mat disparityPart1 = disparityHdr.rowRange(leftFrame.rows / 2, leftFrame.rows);
+
+ gpu::setDevice(0);
- streams[1]->enqueueUpload(leftFrameHdr.rowRange(leftFrame.rows / 2 - 32, leftFrame.rows), d_leftFrames[1]);
- streams[1]->enqueueUpload(rightFrameHdr.rowRange(leftFrame.rows / 2 - 32, leftFrame.rows), d_rightFrames[1]);
- (*d_algs[1])(d_leftFrames[1], d_rightFrames[1], d_disparities[1], *streams[1]);
- streams[1]->enqueueDownload(d_disparities[1].rowRange(32, d_disparities[1].rows), disparityPart1);
++ d_leftFrames[0].upload(leftFrameHdr.rowRange(0, leftFrame.rows / 2 + 32), *streams[0]);
++ d_rightFrames[0].upload(rightFrameHdr.rowRange(0, leftFrame.rows / 2 + 32), *streams[0]);
++ d_algs[0]->compute(d_leftFrames[0], d_rightFrames[0], d_disparities[0], *streams[0]);
++ d_disparities[0].rowRange(0, leftFrame.rows / 2).download(disparityPart0, *streams[0]);
+
+ gpu::setDevice(1);
++ d_leftFrames[1].upload(leftFrameHdr.rowRange(leftFrame.rows / 2 - 32, leftFrame.rows), *streams[1]);
++ d_rightFrames[1].upload(rightFrameHdr.rowRange(leftFrame.rows / 2 - 32, leftFrame.rows), *streams[1]);
++ d_algs[1]->compute(d_leftFrames[1], d_rightFrames[1], d_disparities[1], *streams[1]);
++ d_disparities[1].rowRange(32, d_disparities[1].rows).download(disparityPart1, *streams[1]);
+
+ gpu::setDevice(0);
+ streams[0]->waitForCompletion();
+
+ gpu::setDevice(1);
+ streams[1]->waitForCompletion();
+ }
+
+ ///////////////////////////////////////////////////////////
+ // main
+
int main(int argc, char** argv)
{
- if (argc < 5)
+ if (argc != 3)
{
- printHelp();
+ cerr << "Usage: stereo_multi_gpu <left_video> <right_video>" << endl;
return -1;
}
- int num_devices = getCudaEnabledDeviceCount();
- if (num_devices < 2)
+ const int numDevices = getCudaEnabledDeviceCount();
+ if (numDevices != 2)
{
- std::cout << "Two or more GPUs are required\n";
+ cerr << "Two GPUs are required" << endl;
return -1;
}
- for (int i = 0; i < num_devices; ++i)
- {
- cv::gpu::printShortCudaDeviceInfo(i);
- DeviceInfo dev_info(i);
- if (!dev_info.isCompatible())
+ for (int i = 0; i < numDevices; ++i)
+ {
+ DeviceInfo devInfo(i);
+ if (!devInfo.isCompatible())
{
- std::cout << "GPU module isn't built for GPU #" << i << " ("
- << dev_info.name() << ", CC " << dev_info.majorVersion()
- << dev_info.minorVersion() << "\n";
+ cerr << "GPU module was't built for GPU #" << i << " ("
+ << devInfo.name() << ", CC " << devInfo.majorVersion()
+ << devInfo.minorVersion() << endl;
return -1;
}
+
+ printShortCudaDeviceInfo(i);
}
- // Load input data
- Mat left, right;
- for (int i = 1; i < argc; ++i)
+ VideoCapture leftVideo(argv[1]);
+ VideoCapture rightVideo(argv[2]);
+
+ if (!leftVideo.isOpened())
{
- if (string(argv[i]) == "--left")
- {
- left = imread(argv[++i], cv::IMREAD_GRAYSCALE);
- CV_Assert(!left.empty());
- }
- else if (string(argv[i]) == "--right")
- {
- right = imread(argv[++i], cv::IMREAD_GRAYSCALE);
- CV_Assert(!right.empty());
- }
- else if (string(argv[i]) == "--help")
+ cerr << "Can't open " << argv[1] << " video file" << endl;
+ return -1;
+ }
+
+ if (!rightVideo.isOpened())
+ {
+ cerr << "Can't open " << argv[2] << " video file" << endl;
+ return -1;
+ }
+
+ cout << endl;
+ cout << "This sample demonstrates working on one piece of data using two GPUs." << endl;
+ cout << "It splits input into two parts and processes them separately on different GPUs." << endl;
+ cout << endl;
+
+ Mat leftFrame, rightFrame;
+ CudaMem leftGrayFrame, rightGrayFrame;
+
+ StereoSingleGpu gpu0Alg(0);
+ StereoSingleGpu gpu1Alg(1);
+ StereoMultiGpuThread multiThreadAlg;
+ StereoMultiGpuStream multiStreamAlg;
+
+ Mat disparityGpu0;
+ Mat disparityGpu1;
+ Mat disparityMultiThread;
+ CudaMem disparityMultiStream;
+
+ Mat disparityGpu0Show;
+ Mat disparityGpu1Show;
+ Mat disparityMultiThreadShow;
+ Mat disparityMultiStreamShow;
+
+ TickMeter tm;
+
+ cout << "-------------------------------------------------------------------" << endl;
+ cout << "| Frame | GPU 0 ms | GPU 1 ms | Multi Thread ms | Multi Stream ms |" << endl;
+ cout << "-------------------------------------------------------------------" << endl;
+
+ for (int i = 0;; ++i)
+ {
+ leftVideo >> leftFrame;
+ rightVideo >> rightFrame;
+
+ if (leftFrame.empty() || rightFrame.empty())
+ break;
+
+ if (leftFrame.size() != rightFrame.size())
{
- printHelp();
+ cerr << "Frames have different sizes" << endl;
return -1;
}
- }
- // Split source images for processing on the GPU #0
- setDevice(0);
- d_left[0].upload(left.rowRange(0, left.rows / 2));
- d_right[0].upload(right.rowRange(0, right.rows / 2));
- bm[0] = gpu::createStereoBM();
-
- // Split source images for processing on the GPU #1
- setDevice(1);
- d_left[1].upload(left.rowRange(left.rows / 2, left.rows));
- d_right[1].upload(right.rowRange(right.rows / 2, right.rows));
- bm[1] = gpu::createStereoBM();
-
- // Execute calculation in two threads using two GPUs
- int devices[] = {0, 1};
- tbb::parallel_do(devices, devices + 2, Worker());
-
- // Release the first GPU resources
- setDevice(0);
- imshow("GPU #0 result", Mat(d_result[0]));
- d_left[0].release();
- d_right[0].release();
- d_result[0].release();
- bm[0].release();
-
- // Release the second GPU resources
- setDevice(1);
- imshow("GPU #1 result", Mat(d_result[1]));
- d_left[1].release();
- d_right[1].release();
- d_result[1].release();
- bm[1].release();
-
- waitKey();
- return 0;
- }
+ leftGrayFrame.create(leftFrame.size(), CV_8UC1);
+ rightGrayFrame.create(leftFrame.size(), CV_8UC1);
+ cvtColor(leftFrame, leftGrayFrame.createMatHeader(), COLOR_BGR2GRAY);
+ cvtColor(rightFrame, rightGrayFrame.createMatHeader(), COLOR_BGR2GRAY);
- void Worker::operator()(int device_id) const
- {
- setDevice(device_id);
+ tm.reset(); tm.start();
- gpu0Alg.compute(leftGrayFrame, rightGrayFrame, disparityGpu0);
++ gpu0Alg.compute(leftGrayFrame.createMatHeader(), rightGrayFrame.createMatHeader(),
++ disparityGpu0);
+ tm.stop();
- bm[device_id]->compute(d_left[device_id], d_right[device_id], d_result[device_id]);
+ const double gpu0Time = tm.getTimeMilli();
- std::cout << "GPU #" << device_id << " (" << DeviceInfo().name()
- << "): finished\n";
- }
+ tm.reset(); tm.start();
- gpu1Alg.compute(leftGrayFrame, rightGrayFrame, disparityGpu1);
++ gpu1Alg.compute(leftGrayFrame.createMatHeader(), rightGrayFrame.createMatHeader(),
++ disparityGpu1);
+ tm.stop();
- #endif
+ const double gpu1Time = tm.getTimeMilli();
+
+ tm.reset(); tm.start();
- multiThreadAlg.compute(leftGrayFrame, rightGrayFrame, disparityMultiThread);
++ multiThreadAlg.compute(leftGrayFrame.createMatHeader(), rightGrayFrame.createMatHeader(),
++ disparityMultiThread);
+ tm.stop();
+
+ const double multiThreadTime = tm.getTimeMilli();
+
+ tm.reset(); tm.start();
+ multiStreamAlg.compute(leftGrayFrame, rightGrayFrame, disparityMultiStream);
+ tm.stop();
+
+ const double multiStreamTime = tm.getTimeMilli();
+
+ cout << "| " << setw(5) << i << " | "
+ << setw(8) << setprecision(1) << fixed << gpu0Time << " | "
+ << setw(8) << setprecision(1) << fixed << gpu1Time << " | "
+ << setw(15) << setprecision(1) << fixed << multiThreadTime << " | "
+ << setw(15) << setprecision(1) << fixed << multiStreamTime << " |" << endl;
+
+ resize(disparityGpu0, disparityGpu0Show, Size(1024, 768), 0, 0, INTER_AREA);
+ resize(disparityGpu1, disparityGpu1Show, Size(1024, 768), 0, 0, INTER_AREA);
+ resize(disparityMultiThread, disparityMultiThreadShow, Size(1024, 768), 0, 0, INTER_AREA);
+ resize(disparityMultiStream.createMatHeader(), disparityMultiStreamShow, Size(1024, 768), 0, 0, INTER_AREA);
+
+ imshow("disparityGpu0", disparityGpu0Show);
+ imshow("disparityGpu1", disparityGpu1Show);
+ imshow("disparityMultiThread", disparityMultiThreadShow);
+ imshow("disparityMultiStream", disparityMultiStreamShow);
+
+ const int key = waitKey(30) & 0xff;
+ if (key == 27)
+ break;
+ }
+
+ cout << "-------------------------------------------------------------------" << endl;
+
+ return 0;
+ }
projects / platform / upstream / opencv.git / commitdiff