Merge remote-tracking branch 'origin/2.4' into merge-2.4

Conflicts:
	.gitignore
	modules/contrib/src/detection_based_tracker.cpp
	modules/core/include/opencv2/core/core.hpp
	modules/core/include/opencv2/core/internal.hpp
	modules/core/src/gpumat.cpp
	modules/core/src/opengl.cpp
	modules/gpu/src/cuda/safe_call.hpp
	modules/highgui/src/cap.cpp
	modules/imgproc/include/opencv2/imgproc/imgproc.hpp
	modules/ocl/doc/image_processing.rst
	modules/ocl/include/opencv2/ocl/ocl.hpp
	modules/ocl/perf/perf_haar.cpp
	modules/ocl/src/haar.cpp
	modules/ocl/src/imgproc.cpp
	modules/ocl/src/kmeans.cpp
	modules/ocl/src/svm.cpp
	modules/ocl/test/test_objdetect.cpp
	samples/ocl/adaptive_bilateral_filter.cpp
pull/1781/head
Roman Donchenko 11 years ago
commit 4f1aed98de
  1. 13
      .gitignore
  2. 4
      doc/tutorials/introduction/android_binary_package/O4A_SDK.rst
  3. 27
      modules/contrib/src/detection_based_tracker.cpp
  4. 2
      modules/core/include/opencv2/core.hpp
  5. 16
      modules/core/include/opencv2/core/base.hpp
  6. 6
      modules/core/include/opencv2/core/cuda/common.hpp
  7. 9
      modules/core/include/opencv2/core/private.cuda.hpp
  8. 1
      modules/core/src/copy.cpp
  9. 6
      modules/core/src/opengl.cpp
  10. 7
      modules/core/src/persistence.cpp
  11. 12
      modules/cudaarithm/src/arithm.cpp
  12. 6
      modules/cudalegacy/include/opencv2/cudalegacy/private.hpp
  13. 2
      modules/highgui/src/cap.cpp
  14. 17
      modules/imgproc/doc/filtering.rst
  15. 2
      modules/imgproc/include/opencv2/imgproc.hpp
  16. 85
      modules/imgproc/src/smooth.cpp
  17. 3
      modules/java/CMakeLists.txt
  18. 6
      modules/ocl/doc/camera_calibration_and_3D_reconstruction.rst
  19. 372
      modules/ocl/doc/data_structures.rst
  20. 12
      modules/ocl/doc/image_filtering.rst
  21. 10
      modules/ocl/doc/image_processing.rst
  22. 53
      modules/ocl/doc/introduction.rst
  23. 12
      modules/ocl/doc/ml_machine_learning.rst
  24. 56
      modules/ocl/doc/video_analysis.rst
  25. 21
      modules/ocl/include/opencv2/ocl.hpp
  26. 83
      modules/ocl/include/opencv2/ocl/private/opencl_dumpinfo.hpp
  27. 13
      modules/ocl/perf/main.cpp
  28. 37
      modules/ocl/perf/perf_arithm.cpp
  29. 1
      modules/ocl/perf/perf_haar.cpp
  30. 41
      modules/ocl/perf/perf_imgproc.cpp
  31. 18
      modules/ocl/src/arithm.cpp
  32. 112
      modules/ocl/src/cl_operations.cpp
  33. 30
      modules/ocl/src/filtering.cpp
  34. 32
      modules/ocl/src/imgproc.cpp
  35. 94
      modules/ocl/src/kmeans.cpp
  36. 22
      modules/ocl/src/opencl/bgfg_mog.cl
  37. 63
      modules/ocl/src/opencl/filtering_adaptive_bilateral.cl
  38. 1
      modules/ocl/src/opencl/haarobjectdetect.cl
  39. 2
      modules/ocl/src/opencl/haarobjectdetect_scaled2.cl
  40. 4
      modules/ocl/src/opencl/imgproc_calcHarris.cl
  41. 4
      modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl
  42. 41
      modules/ocl/src/opencl/imgproc_remap.cl
  43. 4
      modules/ocl/src/opencl/imgproc_threshold.cl
  44. 101
      modules/ocl/src/opencl/kmeans_kernel.cl
  45. 16
      modules/ocl/src/opencl/tvl1flow.cl
  46. 9
      modules/ocl/src/safe_call.hpp
  47. 25
      modules/ocl/src/svm.cpp
  48. 12
      modules/ocl/test/main.cpp
  49. 51
      modules/ocl/test/test_arithm.cpp
  50. 4
      modules/ocl/test/test_filters.cpp
  51. 38
      modules/ocl/test/test_imgproc.cpp
  52. 67
      modules/ocl/test/test_kmeans.cpp
  53. 4
      modules/ocl/test/test_ml.cpp
  54. 5
      modules/ocl/test/test_objdetect.cpp
  55. 1
      modules/ocl/test/test_warp.cpp
  56. 8
      modules/world/CMakeLists.txt
  57. 4
      platforms/android/service/engine/AndroidManifest.xml
  58. 28
      platforms/android/service/readme.txt
  59. 31
      samples/ocl/adaptive_bilateral_filter.cpp

13
.gitignore vendored

@ -1,11 +1,10 @@
*.autosave
*.pyc
.DS_Store
refman.rst
OpenCV4Tegra/
tegra/
*.user
.sw[a-z]
.*.swp
tags
.DS_Store
.sw[a-z]
/modules/refman.rst
Thumbs.db
*.autosave
tags
tegra/

@ -51,7 +51,7 @@ The structure of package contents looks as follows:
OpenCV-2.4.7-android-sdk
|_ apk
| |_ OpenCV_2.4.7_binary_pack_armv7a.apk
| |_ OpenCV_2.4.7_Manager_2.13_XXX.apk
| |_ OpenCV_2.4.7_Manager_2.14_XXX.apk
|
|_ doc
|_ samples
@ -295,7 +295,7 @@ Well, running samples from Eclipse is very simple:
.. code-block:: sh
:linenos:
<Android SDK path>/platform-tools/adb install <OpenCV4Android SDK path>/apk/OpenCV_2.4.7_Manager_2.13_armv7a-neon.apk
<Android SDK path>/platform-tools/adb install <OpenCV4Android SDK path>/apk/OpenCV_2.4.7_Manager_2.14_armv7a-neon.apk
.. note:: ``armeabi``, ``armv7a-neon``, ``arm7a-neon-android8``, ``mips`` and ``x86`` stand for
platform targets:

@ -180,37 +180,20 @@ bool cv::DetectionBasedTracker::SeparateDetectionWork::run()
return true;
}
#ifdef __GNUC__
#define CATCH_ALL_AND_LOG(_block) \
do { \
#define CATCH_ALL_AND_LOG(_block) \
do { \
try { \
_block; \
break; \
} \
catch(cv::Exception& e) { \
LOGE0("\n %s: ERROR: OpenCV Exception caught: \n'%s'\n\n", __func__, e.what()); \
LOGE0("\n %s: ERROR: OpenCV Exception caught: \n'%s'\n\n", CV_Func, e.what()); \
} catch(std::exception& e) { \
LOGE0("\n %s: ERROR: Exception caught: \n'%s'\n\n", __func__, e.what()); \
LOGE0("\n %s: ERROR: Exception caught: \n'%s'\n\n", CV_Func, e.what()); \
} catch(...) { \
LOGE0("\n %s: ERROR: UNKNOWN Exception caught\n\n", __func__); \
LOGE0("\n %s: ERROR: UNKNOWN Exception caught\n\n", CV_Func); \
} \
} while(0)
#else
#define CATCH_ALL_AND_LOG(_block) \
do { \
try { \
_block; \
break; \
} \
catch(cv::Exception& e) { \
LOGE0("\n ERROR: OpenCV Exception caught: \n'%s'\n\n", e.what()); \
} catch(std::exception& e) { \
LOGE0("\n ERROR: Exception caught: \n'%s'\n\n", e.what()); \
} catch(...) { \
LOGE0("\n ERROR: UNKNOWN Exception caught\n\n"); \
} \
} while(0)
#endif
void* cv::workcycleObjectDetectorFunction(void* p)
{

@ -90,7 +90,7 @@ public:
int code; ///< error code @see CVStatus
String err; ///< error description
String func; ///< function name. Available only when the compiler supports __func__ macro
String func; ///< function name. Available only when the compiler supports getting it
String file; ///< source file name where the error has occured
int line; ///< line number in the source file where the error has occured
};

@ -222,16 +222,18 @@ enum {
*/
CV_EXPORTS void error(int _code, const String& _err, const char* _func, const char* _file, int _line);
#ifdef __GNUC__
# define CV_Error( code, msg ) cv::error( code, msg, __func__, __FILE__, __LINE__ )
# define CV_Error_( code, args ) cv::error( code, cv::format args, __func__, __FILE__, __LINE__ )
# define CV_Assert( expr ) if(!!(expr)) ; else cv::error( cv::Error::StsAssert, #expr, __func__, __FILE__, __LINE__ )
#if defined __GNUC__
#define CV_Func __func__
#elif defined _MSC_VER
#define CV_Func __FUNCTION__
#else
# define CV_Error( code, msg ) cv::error( code, msg, "", __FILE__, __LINE__ )
# define CV_Error_( code, args ) cv::error( code, cv::format args, "", __FILE__, __LINE__ )
# define CV_Assert( expr ) if(!!(expr)) ; else cv::error( cv::Error::StsAssert, #expr, "", __FILE__, __LINE__ )
#define CV_Func ""
#endif
#define CV_Error( code, msg ) cv::error( code, msg, CV_Func, __FILE__, __LINE__ )
#define CV_Error_( code, args ) cv::error( code, cv::format args, CV_Func, __FILE__, __LINE__ )
#define CV_Assert( expr ) if(!!(expr)) ; else cv::error( cv::Error::StsAssert, #expr, CV_Func, __FILE__, __LINE__ )
#ifdef _DEBUG
# define CV_DbgAssert(expr) CV_Assert(expr)
#else

@ -65,11 +65,7 @@ namespace cv { namespace cuda {
}}
#ifndef cudaSafeCall
#if defined(__GNUC__)
#define cudaSafeCall(expr) cv::cuda::checkCudaError(expr, __FILE__, __LINE__, __func__)
#else /* defined(__CUDACC__) || defined(__MSVC__) */
#define cudaSafeCall(expr) cv::cuda::checkCudaError(expr, __FILE__, __LINE__, "")
#endif
#define cudaSafeCall(expr) cv::cuda::checkCudaError(expr, __FILE__, __LINE__, CV_Func)
#endif
namespace cv { namespace cuda

@ -162,13 +162,8 @@ namespace cv { namespace cuda
};
}}
#if defined(__GNUC__)
#define nppSafeCall(expr) cv::cuda::checkNppError(expr, __FILE__, __LINE__, __func__)
#define cuSafeCall(expr) cv::cuda::checkCudaDriverApiError(expr, __FILE__, __LINE__, __func__)
#else /* defined(__CUDACC__) || defined(__MSVC__) */
#define nppSafeCall(expr) cv::cuda::checkNppError(expr, __FILE__, __LINE__, "")
#define cuSafeCall(expr) cv::cuda::checkCudaDriverApiError(expr, __FILE__, __LINE__, "")
#endif
#define nppSafeCall(expr) cv::cuda::checkNppError(expr, __FILE__, __LINE__, CV_Func)
#define cuSafeCall(expr) cv::cuda::checkCudaDriverApiError(expr, __FILE__, __LINE__, CV_Func)
#endif // HAVE_CUDA

@ -485,6 +485,7 @@ void repeat(InputArray _src, int ny, int nx, OutputArray _dst)
{
Mat src = _src.getMat();
CV_Assert( src.dims <= 2 );
CV_Assert( ny > 0 && nx > 0 );
_dst.create(src.rows*ny, src.cols*nx, src.type());
Mat dst = _dst.getMat();

@ -105,11 +105,7 @@ namespace
#endif
}
#if defined(__GNUC__)
#define CV_CheckGlError() CV_DbgAssert( (checkError(__FILE__, __LINE__, __func__)) )
#else
#define CV_CheckGlError() CV_DbgAssert( (checkError(__FILE__, __LINE__)) )
#endif
#define CV_CheckGlError() CV_DbgAssert( (checkError(__FILE__, __LINE__, CV_Func)) )
} // namespace
#ifdef HAVE_OPENGL

@ -385,13 +385,8 @@ cvCreateMap( int flags, int header_size, int elem_size,
return map;
}
#ifdef __GNUC__
#define CV_PARSE_ERROR( errmsg ) \
icvParseError( fs, __func__, (errmsg), __FILE__, __LINE__ )
#else
#define CV_PARSE_ERROR( errmsg ) \
icvParseError( fs, "", (errmsg), __FILE__, __LINE__ )
#endif
icvParseError( fs, CV_Func, (errmsg), __FILE__, __LINE__ )
static void
icvParseError( CvFileStorage* fs, const char* func_name,

@ -113,11 +113,7 @@ namespace
}
}
#if defined(__GNUC__)
#define cublasSafeCall(expr) ___cublasSafeCall(expr, __FILE__, __LINE__, __func__)
#else /* defined(__CUDACC__) || defined(__MSVC__) */
#define cublasSafeCall(expr) ___cublasSafeCall(expr, __FILE__, __LINE__, "")
#endif
#define cublasSafeCall(expr) ___cublasSafeCall(expr, __FILE__, __LINE__, CV_Func)
#endif // HAVE_CUBLAS
#ifdef HAVE_CUFFT
@ -151,11 +147,7 @@ namespace
}
}
#if defined(__GNUC__)
#define cufftSafeCall(expr) ___cufftSafeCall(expr, __FILE__, __LINE__, __func__)
#else /* defined(__CUDACC__) || defined(__MSVC__) */
#define cufftSafeCall(expr) ___cufftSafeCall(expr, __FILE__, __LINE__, "")
#endif
#define cufftSafeCall(expr) ___cufftSafeCall(expr, __FILE__, __LINE__, CV_Func)
#endif

@ -87,10 +87,6 @@ namespace cv { namespace cuda
}
}}
#if defined(__GNUC__)
#define ncvSafeCall(expr) cv::cuda::checkNcvError(expr, __FILE__, __LINE__, __func__)
#else /* defined(__CUDACC__) || defined(__MSVC__) */
#define ncvSafeCall(expr) cv::cuda::checkNcvError(expr, __FILE__, __LINE__, "")
#endif
#define ncvSafeCall(expr) cv::cuda::checkNcvError(expr, __FILE__, __LINE__, CV_Func)
#endif // __OPENCV_CORE_CUDALEGACY_PRIVATE_HPP__

@ -524,7 +524,7 @@ bool VideoCapture::retrieve(Mat& image, int channel)
return false;
}
if(_img->origin == IPL_ORIGIN_TL)
image = cv::cvarrToMat(_img);
cv::cvarrToMat(_img).copyTo(image);
else
{
Mat temp = cv::cvarrToMat(_img);

@ -416,24 +416,23 @@ adaptiveBilateralFilter
-----------------------
Applies the adaptive bilateral filter to an image.
.. ocv:function:: void adaptiveBilateralFilter( InputArray src, OutputArray dst, Size ksize, double sigmaSpace, Point anchor=Point(-1, -1), int borderType=BORDER_DEFAULT )
.. ocv:function:: void adaptiveBilateralFilter( InputArray src, OutputArray dst, Size ksize, double sigmaSpace, double maxSigmaColor = 20.0, Point anchor=Point(-1, -1), int borderType=BORDER_DEFAULT )
.. ocv:pyfunction:: cv2.adaptiveBilateralFilter(src, ksize, sigmaSpace[, dst[, anchor[, borderType]]]) -> dst
:param src: Source 8-bit, 1-channel or 3-channel image.
:param dst: Destination image of the same size and type as ``src`` .
:param src: The source image
:param ksize: filter kernel size.
:param dst: The destination image; will have the same size and the same type as src
:param sigmaSpace: Filter sigma in the coordinate space. It has similar meaning with ``sigmaSpace`` in ``bilateralFilter``.
:param ksize: The kernel size. This is the neighborhood where the local variance will be calculated, and where pixels will contribute (in a weighted manner).
:param anchor: anchor point; default value ``Point(-1,-1)`` means that the anchor is at the kernel center. Only default value is supported now.
:param sigmaSpace: Filter sigma in the coordinate space. Larger value of the parameter means that farther pixels will influence each other (as long as their colors are close enough; see sigmaColor). Then d>0, it specifies the neighborhood size regardless of sigmaSpace, otherwise d is proportional to sigmaSpace.
:param borderType: border mode used to extrapolate pixels outside of the image.
:param maxSigmaColor: Maximum allowed sigma color (will clamp the value calculated in the ksize neighborhood. Larger value of the parameter means that more dissimilar pixels will influence each other (as long as their colors are close enough; see sigmaColor). Then d>0, it specifies the neighborhood size regardless of sigmaSpace, otherwise d is proportional to sigmaSpace.
The function applies adaptive bilateral filtering to the input image. This filter is similar to ``bilateralFilter``, in that dissimilarity from and distance to the center pixel is punished. Instead of using ``sigmaColor``, we employ the variance of pixel values in the neighbourhood.
:param borderType: Pixel extrapolation method.
A main part of our strategy will be to load each raw pixel once, and reuse it to calculate all pixels in the output (filtered) image that need this pixel value. The math of the filter is that of the usual bilateral filter, except that the sigma color is calculated in the neighborhood, and clamped by the optional input value.
blur

@ -1067,7 +1067,7 @@ CV_EXPORTS_W void bilateralFilter( InputArray src, OutputArray dst, int d,
//! smooths the image using adaptive bilateral filter
CV_EXPORTS_W void adaptiveBilateralFilter( InputArray src, OutputArray dst, Size ksize,
double sigmaSpace, Point anchor=Point(-1, -1),
double sigmaSpace, double maxSigmaColor = 20.0, Point anchor=Point(-1, -1),
int borderType=BORDER_DEFAULT );
//! smooths the image using the box filter. Each pixel is processed in O(1) time

@ -2287,15 +2287,24 @@ void cv::bilateralFilter( InputArray _src, OutputArray _dst, int d,
namespace cv
{
#define CALCVAR 1
#define FIXED_WEIGHT 0
#ifndef ABF_CALCVAR
#define ABF_CALCVAR 1
#endif
#ifndef ABF_FIXED_WEIGHT
#define ABF_FIXED_WEIGHT 0
#endif
#ifndef ABF_GAUSSIAN
#define ABF_GAUSSIAN 1
#endif
class adaptiveBilateralFilter_8u_Invoker :
public ParallelLoopBody
{
public:
adaptiveBilateralFilter_8u_Invoker(Mat& _dest, const Mat& _temp, Size _ksize, double _sigma_space, Point _anchor) :
temp(&_temp), dest(&_dest), ksize(_ksize), sigma_space(_sigma_space), anchor(_anchor)
adaptiveBilateralFilter_8u_Invoker(Mat& _dest, const Mat& _temp, Size _ksize, double _sigma_space, double _maxSigmaColor, Point _anchor) :
temp(&_temp), dest(&_dest), ksize(_ksize), sigma_space(_sigma_space), maxSigma_Color(_maxSigmaColor), anchor(_anchor)
{
if( sigma_space <= 0 )
sigma_space = 1;
@ -2308,7 +2317,11 @@ public:
for(int y=-h; y<=h; y++)
for(int x=-w; x<=w; x++)
{
#if ABF_GAUSSIAN
space_weight[idx++] = (float)exp ( -0.5*(x * x + y * y)/sigma2);
#else
space_weight[idx++] = (float)(sigma2 / (sigma2 + x * x + y * y));
#endif
}
}
virtual void operator()(const Range& range) const
@ -2344,7 +2357,7 @@ public:
int startLMJ = 0;
int endLMJ = ksize.width - 1;
int howManyAll = (anX *2 +1)*(ksize.width );
#if CALCVAR
#if ABF_CALCVAR
for(int x = startLMJ; x< endLMJ; x++)
{
tptr = temp->ptr(startY + x) +j;
@ -2356,8 +2369,14 @@ public:
}
}
var = ( (sumValSqr * howManyAll)- sumVal * sumVal ) / ( (float)(howManyAll*howManyAll));
if(var < 0.01)
var = 0.01f;
else if(var > (float)(maxSigma_Color*maxSigma_Color) )
var = (float)(maxSigma_Color*maxSigma_Color) ;
#else
var = 900.0;
var = maxSigmaColor*maxSigmaColor;
#endif
startLMJ = 0;
endLMJ = ksize.width;
@ -2368,13 +2387,18 @@ public:
tptr = temp->ptr(startY + x) +j;
for(int y=-anX; y<=anX; y++)
{
#if FIXED_WEIGHT
#if ABF_FIXED_WEIGHT
weight = 1.0;
#else
currVal = tptr[cn*(y+anX)];
currWRTCenter = currVal - currValCenter;
weight = var / ( var + (currWRTCenter * currWRTCenter) ) * space_weight[x*ksize.width+y+anX];;
#if ABF_GAUSSIAN
weight = exp ( -0.5f * currWRTCenter * currWRTCenter/var ) * space_weight[x*ksize.width+y+anX];
#else
weight = var / ( var + (currWRTCenter * currWRTCenter) ) * space_weight[x*ksize.width+y+anX];
#endif
#endif
tmpSum += ((float)tptr[cn*(y+anX)] * weight);
totalWeight += weight;
@ -2409,7 +2433,8 @@ public:
int startLMJ = 0;
int endLMJ = ksize.width - 1;
int howManyAll = (anX *2 +1)*(ksize.width);
#if CALCVAR
#if ABF_CALCVAR
float max_var = (float)( maxSigma_Color*maxSigma_Color);
for(int x = startLMJ; x< endLMJ; x++)
{
tptr = temp->ptr(startY + x) +j;
@ -2424,11 +2449,27 @@ public:
sumValSqr_r += (currVal_r *currVal_r);
}
}
var_b = ( (sumValSqr_b * howManyAll)- sumVal_b * sumVal_b ) / ( (float)(howManyAll*howManyAll));
var_g = ( (sumValSqr_g * howManyAll)- sumVal_g * sumVal_g ) / ( (float)(howManyAll*howManyAll));
var_r = ( (sumValSqr_r * howManyAll)- sumVal_r * sumVal_r ) / ( (float)(howManyAll*howManyAll));
var_b = ( (sumValSqr_b * howManyAll)- sumVal_b * sumVal_b ) / ( (float)(howManyAll*howManyAll));
var_g = ( (sumValSqr_g * howManyAll)- sumVal_g * sumVal_g ) / ( (float)(howManyAll*howManyAll));
var_r = ( (sumValSqr_r * howManyAll)- sumVal_r * sumVal_r ) / ( (float)(howManyAll*howManyAll));
if(var_b < 0.01)
var_b = 0.01f;
else if(var_b > max_var )
var_b = (float)(max_var) ;
if(var_g < 0.01)
var_g = 0.01f;
else if(var_g > max_var )
var_g = (float)(max_var) ;
if(var_r < 0.01)
var_r = 0.01f;
else if(var_r > max_var )
var_r = (float)(max_var) ;
#else
var_b = 900.0; var_g = 900.0;var_r = 900.0;
var_b = maxSigma_Color*maxSigma_Color; var_g = maxSigma_Color*maxSigma_Color; var_r = maxSigma_Color*maxSigma_Color;
#endif
startLMJ = 0;
endLMJ = ksize.width;
@ -2439,7 +2480,7 @@ public:
tptr = temp->ptr(startY + x) +j;
for(int y=-anX; y<=anX; y++)
{
#if FIXED_WEIGHT
#if ABF_FIXED_WEIGHT
weight_b = 1.0;
weight_g = 1.0;
weight_r = 1.0;
@ -2450,9 +2491,16 @@ public:
currWRTCenter_r = currVal_r - currValCenter_r;
float cur_spw = space_weight[x*ksize.width+y+anX];
#if ABF_GAUSSIAN
weight_b = exp( -0.5f * currWRTCenter_b * currWRTCenter_b/ var_b ) * cur_spw;
weight_g = exp( -0.5f * currWRTCenter_g * currWRTCenter_g/ var_g ) * cur_spw;
weight_r = exp( -0.5f * currWRTCenter_r * currWRTCenter_r/ var_r ) * cur_spw;
#else
weight_b = var_b / ( var_b + (currWRTCenter_b * currWRTCenter_b) ) * cur_spw;
weight_g = var_g / ( var_g + (currWRTCenter_g * currWRTCenter_g) ) * cur_spw;
weight_r = var_r / ( var_r + (currWRTCenter_r * currWRTCenter_r) ) * cur_spw;
#endif
#endif
tmpSum_b += ((float)tptr[cn*(y+anX)] * weight_b);
tmpSum_g += ((float)tptr[cn*(y+anX)+1] * weight_g);
@ -2476,10 +2524,11 @@ private:
Mat *dest;
Size ksize;
double sigma_space;
double maxSigma_Color;
Point anchor;
std::vector<float> space_weight;
};
static void adaptiveBilateralFilter_8u( const Mat& src, Mat& dst, Size ksize, double sigmaSpace, Point anchor, int borderType )
static void adaptiveBilateralFilter_8u( const Mat& src, Mat& dst, Size ksize, double sigmaSpace, double maxSigmaColor, Point anchor, int borderType )
{
Size size = src.size();
@ -2489,12 +2538,12 @@ static void adaptiveBilateralFilter_8u( const Mat& src, Mat& dst, Size ksize, do
Mat temp;
copyMakeBorder(src, temp, anchor.x, anchor.y, anchor.x, anchor.y, borderType);
adaptiveBilateralFilter_8u_Invoker body(dst, temp, ksize, sigmaSpace, anchor);
adaptiveBilateralFilter_8u_Invoker body(dst, temp, ksize, sigmaSpace, maxSigmaColor, anchor);
parallel_for_(Range(0, size.height), body, dst.total()/(double)(1<<16));
}
}
void cv::adaptiveBilateralFilter( InputArray _src, OutputArray _dst, Size ksize,
double sigmaSpace, Point anchor, int borderType )
double sigmaSpace, double maxSigmaColor, Point anchor, int borderType )
{
Mat src = _src.getMat();
_dst.create(src.size(), src.type());
@ -2504,7 +2553,7 @@ void cv::adaptiveBilateralFilter( InputArray _src, OutputArray _dst, Size ksize,
anchor = normalizeAnchor(anchor,ksize);
if( src.depth() == CV_8U )
adaptiveBilateralFilter_8u( src, dst, ksize, sigmaSpace, anchor, borderType );
adaptiveBilateralFilter_8u( src, dst, ksize, sigmaSpace, maxSigmaColor, anchor, borderType );
else
CV_Error( CV_StsUnsupportedFormat,
"Adaptive Bilateral filtering is only implemented for 8u images" );

@ -134,9 +134,10 @@ endforeach()
set(step2_depends ${step1_depends} ${scripts_gen_javadoc} ${scripts_rst_parser} ${javadoc_rst_sources} ${generated_java_sources} ${handwrittren_java_sources})
string(REPLACE ";" "," OPENCV_JAVA_MODULES_STR "${OPENCV_JAVA_MODULES}")
add_custom_command(OUTPUT ${documented_java_files}
COMMAND ${PYTHON_EXECUTABLE} "${scripts_gen_javadoc}" --modules ${OPENCV_JAVA_MODULES_STR} "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/java" "${CMAKE_CURRENT_BINARY_DIR}" 2>"${CMAKE_CURRENT_BINARY_DIR}/get_javadoc_errors.log"
COMMAND ${PYTHON_EXECUTABLE} "${scripts_gen_javadoc}" --modules ${OPENCV_JAVA_MODULES_STR} "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/java" "${CMAKE_CURRENT_BINARY_DIR}" 2> "${CMAKE_CURRENT_BINARY_DIR}/get_javadoc_errors.log"
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
DEPENDS ${step2_depends}
VERBATIM
)
# step 3: copy files to destination

@ -86,8 +86,6 @@ Enables the stereo correspondence operator that finds the disparity for the spec
:param disparity: Output disparity map. It is a ``CV_8UC1`` image with the same size as the input images.
:param stream: Stream for the asynchronous version.
ocl::StereoBM_OCL::checkIfGpuCallReasonable
-----------------------------------------------
@ -218,8 +216,6 @@ Enables the stereo correspondence operator that finds the disparity for the spec
:param disparity: Output disparity map. If ``disparity`` is empty, the output type is ``CV_16SC1`` . Otherwise, the type is retained.
:param stream: Stream for the asynchronous version.
ocl::StereoConstantSpaceBP
------------------------------
.. ocv:class:: ocl::StereoConstantSpaceBP
@ -330,5 +326,3 @@ Enables the stereo correspondence operator that finds the disparity for the spec
:param right: Right image with the same size and the same type as the left one.
:param disparity: Output disparity map. If ``disparity`` is empty, the output type is ``CV_16SC1`` . Otherwise, the output type is ``disparity.type()`` .
:param stream: Stream for the asynchronous version.

@ -5,185 +5,193 @@ Data Structures
OpenCV C++ 1-D or 2-D dense array class ::
class CV_EXPORTS oclMat
{
public:
//! default constructor
oclMat();
//! constructs oclMatrix of the specified size and type (_type is CV_8UC1, CV_64FC3, CV_32SC(12) etc.)
oclMat(int rows, int cols, int type);
oclMat(Size size, int type);
//! constucts oclMatrix and fills it with the specified value _s.
oclMat(int rows, int cols, int type, const Scalar &s);
oclMat(Size size, int type, const Scalar &s);
//! copy constructor
oclMat(const oclMat &m);
//! constructor for oclMatrix headers pointing to user-allocated data
oclMat(int rows, int cols, int type, void *data, size_t step = Mat::AUTO_STEP);
oclMat(Size size, int type, void *data, size_t step = Mat::AUTO_STEP);
//! creates a matrix header for a part of the bigger matrix
oclMat(const oclMat &m, const Range &rowRange, const Range &colRange);
oclMat(const oclMat &m, const Rect &roi);
//! builds oclMat from Mat. Perfom blocking upload to device.
explicit oclMat (const Mat &m);
//! destructor - calls release()
~oclMat();
//! assignment operators
oclMat &operator = (const oclMat &m);
//! assignment operator. Perfom blocking upload to device.
oclMat &operator = (const Mat &m);
//! pefroms blocking upload data to oclMat.
void upload(const cv::Mat &m);
//! downloads data from device to host memory. Blocking calls.
operator Mat() const;
void download(cv::Mat &m) const;
//! returns a new oclMatrix header for the specified row
oclMat row(int y) const;
//! returns a new oclMatrix header for the specified column
oclMat col(int x) const;
//! ... for the specified row span
oclMat rowRange(int startrow, int endrow) const;
oclMat rowRange(const Range &r) const;
//! ... for the specified column span
oclMat colRange(int startcol, int endcol) const;
oclMat colRange(const Range &r) const;
//! returns deep copy of the oclMatrix, i.e. the data is copied
oclMat clone() const;
//! copies the oclMatrix content to "m".
// It calls m.create(this->size(), this->type()).
// It supports any data type
void copyTo( oclMat &m ) const;
//! copies those oclMatrix elements to "m" that are marked with non-zero mask elements.
//It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
void copyTo( oclMat &m, const oclMat &mask ) const;
//! converts oclMatrix to another datatype with optional scalng. See cvConvertScale.
//It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
void convertTo( oclMat &m, int rtype, double alpha = 1, double beta = 0 ) const;
void assignTo( oclMat &m, int type = -1 ) const;
//! sets every oclMatrix element to s
//It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
oclMat &operator = (const Scalar &s);
//! sets some of the oclMatrix elements to s, according to the mask
//It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
oclMat &setTo(const Scalar &s, const oclMat &mask = oclMat());
//! creates alternative oclMatrix header for the same data, with different
// number of channels and/or different number of rows. see cvReshape.
oclMat reshape(int cn, int rows = 0) const;
//! allocates new oclMatrix data unless the oclMatrix already has specified size and type.
// previous data is unreferenced if needed.
void create(int rows, int cols, int type);
void create(Size size, int type);
//! decreases reference counter;
// deallocate the data when reference counter reaches 0.
void release();
//! swaps with other smart pointer
void swap(oclMat &mat);
//! locates oclMatrix header within a parent oclMatrix. See below
void locateROI( Size &wholeSize, Point &ofs ) const;
//! moves/resizes the current oclMatrix ROI inside the parent oclMatrix.
oclMat &adjustROI( int dtop, int dbottom, int dleft, int dright );
//! extracts a rectangular sub-oclMatrix
// (this is a generalized form of row, rowRange etc.)
oclMat operator()( Range rowRange, Range colRange ) const;
oclMat operator()( const Rect &roi ) const;
//! returns true if the oclMatrix data is continuous
// (i.e. when there are no gaps between successive rows).
// similar to CV_IS_oclMat_CONT(cvoclMat->type)
bool isContinuous() const;
//! returns element size in bytes,
// similar to CV_ELEM_SIZE(cvMat->type)
size_t elemSize() const;
//! returns the size of element channel in bytes.
size_t elemSize1() const;
//! returns element type, similar to CV_MAT_TYPE(cvMat->type)
int type() const;
//! returns element type, i.e. 8UC3 returns 8UC4 because in ocl
//! 3 channels element actually use 4 channel space
int ocltype() const;
//! returns element type, similar to CV_MAT_DEPTH(cvMat->type)
int depth() const;
//! returns element type, similar to CV_MAT_CN(cvMat->type)
int channels() const;
//! returns element type, return 4 for 3 channels element,
//!becuase 3 channels element actually use 4 channel space
int oclchannels() const;
//! returns step/elemSize1()
size_t step1() const;
//! returns oclMatrix size:
// width == number of columns, height == number of rows
Size size() const;
//! returns true if oclMatrix data is NULL
bool empty() const;
//! returns pointer to y-th row
uchar *ptr(int y = 0);
const uchar *ptr(int y = 0) const;
//! template version of the above method
template<typename _Tp> _Tp *ptr(int y = 0);
template<typename _Tp> const _Tp *ptr(int y = 0) const;
//! matrix transposition
oclMat t() const;
/*! includes several bit-fields:
- the magic signature
- continuity flag
- depth
- number of channels
*/
int flags;
//! the number of rows and columns
int rows, cols;
//! a distance between successive rows in bytes; includes the gap if any
size_t step;
//! pointer to the data(OCL memory object)
uchar *data;
//! pointer to the reference counter;
// when oclMatrix points to user-allocated data, the pointer is NULL
int *refcount;
//! helper fields used in locateROI and adjustROI
//datastart and dataend are not used in current version
uchar *datastart;
uchar *dataend;
//! OpenCL context associated with the oclMat object.
Context *clCxt;
//add offset for handle ROI, calculated in byte
int offset;
//add wholerows and wholecols for the whole matrix, datastart and dataend are no longer used
int wholerows;
int wholecols;
};
Basically speaking, the oclMat is the mirror of Mat with the extension of ocl feature, the members have the same meaning and useage of Mat except following:
datastart and dataend are replaced with wholerows and wholecols
add clCxt for oclMat
Only basic flags are supported in oclMat(i.e. depth number of channels)
All the 3-channel matrix(i.e. RGB image) are represented by 4-channel matrix in oclMat. It means 3-channel image have 4-channel space with the last channel unused. We provide a transparent interface to handle the difference between OpenCV Mat and oclMat.
For example: If a oclMat has 3 channels, channels() returns 3 and oclchannels() returns 4
class CV_EXPORTS oclMat
{
public:
//! default constructor
oclMat();
//! constructs oclMatrix of the specified size and type (_type is CV_8UC1, CV_64FC3, CV_32SC(12) etc.)
oclMat(int rows, int cols, int type);
oclMat(Size size, int type);
//! constucts oclMatrix and fills it with the specified value _s.
oclMat(int rows, int cols, int type, const Scalar &s);
oclMat(Size size, int type, const Scalar &s);
//! copy constructor
oclMat(const oclMat &m);
//! constructor for oclMatrix headers pointing to user-allocated data
oclMat(int rows, int cols, int type, void *data, size_t step = Mat::AUTO_STEP);
oclMat(Size size, int type, void *data, size_t step = Mat::AUTO_STEP);
//! creates a matrix header for a part of the bigger matrix
oclMat(const oclMat &m, const Range &rowRange, const Range &colRange);
oclMat(const oclMat &m, const Rect &roi);
//! builds oclMat from Mat. Perfom blocking upload to device.
explicit oclMat (const Mat &m);
//! destructor - calls release()
~oclMat();
//! assignment operators
oclMat &operator = (const oclMat &m);
//! assignment operator. Perfom blocking upload to device.
oclMat &operator = (const Mat &m);
oclMat &operator = (const oclMatExpr& expr);
//! pefroms blocking upload data to oclMat.
void upload(const cv::Mat &m);
//! downloads data from device to host memory. Blocking calls.
operator Mat() const;
void download(cv::Mat &m) const;
//! convert to _InputArray
operator _InputArray();
//! convert to _OutputArray
operator _OutputArray();
//! returns a new oclMatrix header for the specified row
oclMat row(int y) const;
//! returns a new oclMatrix header for the specified column
oclMat col(int x) const;
//! ... for the specified row span
oclMat rowRange(int startrow, int endrow) const;
oclMat rowRange(const Range &r) const;
//! ... for the specified column span
oclMat colRange(int startcol, int endcol) const;
oclMat colRange(const Range &r) const;
//! returns deep copy of the oclMatrix, i.e. the data is copied
oclMat clone() const;
//! copies those oclMatrix elements to "m" that are marked with non-zero mask elements.
// It calls m.create(this->size(), this->type()).
// It supports any data type
void copyTo( oclMat &m, const oclMat &mask = oclMat()) const;
//! converts oclMatrix to another datatype with optional scalng. See cvConvertScale.
void convertTo( oclMat &m, int rtype, double alpha = 1, double beta = 0 ) const;
void assignTo( oclMat &m, int type = -1 ) const;
//! sets every oclMatrix element to s
oclMat& operator = (const Scalar &s);
//! sets some of the oclMatrix elements to s, according to the mask
oclMat& setTo(const Scalar &s, const oclMat &mask = oclMat());
//! creates alternative oclMatrix header for the same data, with different
// number of channels and/or different number of rows. see cvReshape.
oclMat reshape(int cn, int rows = 0) const;
//! allocates new oclMatrix data unless the oclMatrix already has specified size and type.
// previous data is unreferenced if needed.
void create(int rows, int cols, int type);
void create(Size size, int type);
//! allocates new oclMatrix with specified device memory type.
void createEx(int rows, int cols, int type, DevMemRW rw_type, DevMemType mem_type);
void createEx(Size size, int type, DevMemRW rw_type, DevMemType mem_type);
//! decreases reference counter;
// deallocate the data when reference counter reaches 0.
void release();
//! swaps with other smart pointer
void swap(oclMat &mat);
//! locates oclMatrix header within a parent oclMatrix. See below
void locateROI( Size &wholeSize, Point &ofs ) const;
//! moves/resizes the current oclMatrix ROI inside the parent oclMatrix.
oclMat& adjustROI( int dtop, int dbottom, int dleft, int dright );
//! extracts a rectangular sub-oclMatrix
// (this is a generalized form of row, rowRange etc.)
oclMat operator()( Range rowRange, Range colRange ) const;
oclMat operator()( const Rect &roi ) const;
oclMat& operator+=( const oclMat& m );
oclMat& operator-=( const oclMat& m );
oclMat& operator*=( const oclMat& m );
oclMat& operator/=( const oclMat& m );
//! returns true if the oclMatrix data is continuous
// (i.e. when there are no gaps between successive rows).
// similar to CV_IS_oclMat_CONT(cvoclMat->type)
bool isContinuous() const;
//! returns element size in bytes,
// similar to CV_ELEM_SIZE(cvMat->type)
size_t elemSize() const;
//! returns the size of element channel in bytes.
size_t elemSize1() const;
//! returns element type, similar to CV_MAT_TYPE(cvMat->type)
int type() const;
//! returns element type, i.e. 8UC3 returns 8UC4 because in ocl
//! 3 channels element actually use 4 channel space
int ocltype() const;
//! returns element type, similar to CV_MAT_DEPTH(cvMat->type)
int depth() const;
//! returns element type, similar to CV_MAT_CN(cvMat->type)
int channels() const;
//! returns element type, return 4 for 3 channels element,
//!becuase 3 channels element actually use 4 channel space
int oclchannels() const;
//! returns step/elemSize1()
size_t step1() const;
//! returns oclMatrix size:
// width == number of columns, height == number of rows
Size size() const;
//! returns true if oclMatrix data is NULL
bool empty() const;
//! returns pointer to y-th row
uchar* ptr(int y = 0);
const uchar *ptr(int y = 0) const;
//! template version of the above method
template<typename _Tp> _Tp *ptr(int y = 0);
template<typename _Tp> const _Tp *ptr(int y = 0) const;
//! matrix transposition
oclMat t() const;
/*! includes several bit-fields:
- the magic signature
- continuity flag
- depth
- number of channels
*/
int flags;
//! the number of rows and columns
int rows, cols;
//! a distance between successive rows in bytes; includes the gap if any
size_t step;
//! pointer to the data(OCL memory object)
uchar *data;
//! pointer to the reference counter;
// when oclMatrix points to user-allocated data, the pointer is NULL
int *refcount;
//! helper fields used in locateROI and adjustROI
//datastart and dataend are not used in current version
uchar *datastart;
uchar *dataend;
//! OpenCL context associated with the oclMat object.
Context *clCxt;
//add offset for handle ROI, calculated in byte
int offset;
//add wholerows and wholecols for the whole matrix, datastart and dataend are no longer used
int wholerows;
int wholecols;
};
Basically speaking, the ``oclMat`` is the mirror of ``Mat`` with the extension of OCL feature, the members have the same meaning and useage of ``Mat`` except following:
* ``datastart`` and ``dataend`` are replaced with ``wholerows`` and ``wholecols``
* Only basic flags are supported in ``oclMat`` (i.e. depth number of channels)
* All the 3-channel matrix (i.e. RGB image) are represented by 4-channel matrix in ``oclMat``. It means 3-channel image have 4-channel space with the last channel unused. We provide a transparent interface to handle the difference between OpenCV ``Mat`` and ``oclMat``.
For example: If a ``oclMat`` has 3 channels, ``channels()`` returns 3 and ``oclchannels()`` returns 4

@ -532,23 +532,21 @@ ocl::adaptiveBilateralFilter
--------------------------------
Returns void
.. ocv:function:: void ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, Point anchor = Point(-1, -1), int borderType=BORDER_DEFAULT)
.. ocv:function:: void ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, double maxSigmaColor = 20.0, Point anchor = Point(-1, -1), int borderType=BORDER_DEFAULT)
:param src: The source image
:param dst: The destination image; will have the same size and the same type as src
:param ksize: The kernel size
:param ksize: The kernel size. This is the neighborhood where the local variance will be calculated, and where pixels will contribute (in a weighted manner).
:param sigmaSpace: Filter sigma in the coordinate space. Larger value of the parameter means that farther pixels will influence each other (as long as their colors are close enough; see sigmaColor). Then d>0, it specifies the neighborhood size regardless of sigmaSpace, otherwise d is proportional to sigmaSpace.
:param borderType: Pixel extrapolation method.
A main part of our strategy will be to load each raw pixel once, and reuse it to calculate all pixels in the output (filtered) image that need this pixel value.
:param maxSigmaColor: Maximum allowed sigma color (will clamp the value calculated in the ksize neighborhood. Larger value of the parameter means that more dissimilar pixels will influence each other (as long as their colors are close enough; see sigmaColor). Then d>0, it specifies the neighborhood size regardless of sigmaSpace, otherwise d is proportional to sigmaSpace.
.. math::
:param borderType: Pixel extrapolation method.
\emph{O}_i = \frac{1}{W_i}\sum\limits_{j\in{N(i)}}{\frac{1}{1+\frac{(V_i-V_j)^2}{\sigma_{N{'}(i)}^2}}*\frac{1}{1+\frac{d(i,j)^2}{\sum^2}}}V_j
A main part of our strategy will be to load each raw pixel once, and reuse it to calculate all pixels in the output (filtered) image that need this pixel value. The math of the filter is that of the usual bilateral filter, except that the sigma color is calculated in the neighborhood, and clamped by the optional input value.
Local memory organization

@ -146,7 +146,7 @@ Returns void
.. ocv:function:: void ocl::remap(const oclMat &src, oclMat &dst, oclMat &map1, oclMat &map2, int interpolation, int bordertype, const Scalar &value = Scalar())
:param src: Source image. Only CV_8UC1 and CV_32FC1 images are supported now.
:param src: Source image.
:param dst: Destination image containing cornerness values. It has the same size as src and CV_32FC1 type.
@ -156,11 +156,11 @@ Returns void
:param interpolation: The interpolation method
:param bordertype: Pixel extrapolation method. Only BORDER_CONSTANT are supported now.
:param bordertype: Pixel extrapolation method.
:param value: The border value if borderType==BORDER CONSTANT
The function remap transforms the source image using the specified map: dst (x ,y) = src (map1(x , y) , map2(x , y)) where values of pixels with non-integer coordinates are computed using one of available interpolation methods. map1 and map2 can be encoded as separate floating-point maps in map1 and map2 respectively, or interleaved floating-point maps of (x,y) in map1. Supports CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1 , CV_32FC3 and CV_32FC4 data types.
The function remap transforms the source image using the specified map: dst (x ,y) = src (map1(x , y) , map2(x , y)) where values of pixels with non-integer coordinates are computed using one of available interpolation methods. map1 and map2 can be encoded as separate floating-point maps in map1 and map2 respectively, or interleaved floating-point maps of (x,y) in map1.
ocl::resize
------------------
@ -222,7 +222,7 @@ ocl::cvtColor
------------------
Returns void
.. ocv:function:: void ocl::cvtColor(const oclMat &src, oclMat &dst, int code , int dcn = 0)
.. ocv:function:: void ocl::cvtColor(const oclMat &src, oclMat &dst, int code, int dcn = 0)
:param src: Source image.
@ -250,7 +250,7 @@ Returns Threshold value
:param type: Thresholding type
The function applies fixed-level thresholding to a single-channel array. The function is typically used to get a bi-level (binary) image out of a grayscale image or for removing a noise, i.e. filtering out pixels with too small or too large values. There are several types of thresholding that the function supports that are determined by thresholdType. Supports only CV_32FC1 and CV_8UC1 data type.
The function applies fixed-level thresholding to a single-channel array. The function is typically used to get a bi-level (binary) image out of a grayscale image or for removing a noise, i.e. filtering out pixels with too small or too large values. There are several types of thresholding that the function supports that are determined by thresholdType.
ocl::buildWarpPlaneMaps
-----------------------

@ -6,53 +6,68 @@ OpenCL Module Introduction
General Information
-------------------
The OpenCV OCL module contains a set of classes and functions that implement and accelerate select openCV functionality on OpenCL compatible devices. OpenCL is a Khronos standard, implemented by a variety of devices (CPUs, GPUs, FPGAs, ARM), abstracting the exact hardware details, while enabling vendors to provide native implementation for maximal acceleration on their hardware. The standard enjoys wide industry support, and the end user of the module will enjoy the data parallelism benefits that the specific platform/hardware may be capable of, in a platform/hardware independent manner.
The OpenCV OCL module contains a set of classes and functions that implement and accelerate OpenCV functionality on OpenCL compatible devices. OpenCL is a Khronos standard, implemented by a variety of devices (CPUs, GPUs, FPGAs, ARM), abstracting the exact hardware details, while enabling vendors to provide native implementation for maximal acceleration on their hardware. The standard enjoys wide industry support, and the end user of the module will enjoy the data parallelism benefits that the specific platform/hardware may be capable of, in a platform/hardware independent manner.
While in the future we hope to validate (and enable) the OCL module in all OpenCL capable devices, we currently develop and test on GPU devices only. This includes both discrete GPUs (NVidia, AMD), as well as integrated chips(AMD APU and intel HD devices). Performance of any particular algorithm will depend on the particular platform characteristics and capabilities. However, currently (as of 2.4.4), accuracy and mathematical correctness has been verified to be identical to that of the pure CPU implementation on all tested GPU devices and platforms (both windows and linux).
While in the future we hope to validate (and enable) the OCL module in all OpenCL capable devices, we currently develop and test on GPU devices only. This includes both discrete GPUs (NVidia, AMD), as well as integrated chips (AMD APU and Intel HD devices). Performance of any particular algorithm will depend on the particular platform characteristics and capabilities. However, currently, accuracy and mathematical correctness has been verified to be identical to that of the pure CPU implementation on all tested GPU devices and platforms (both Windows and Linux).
The OpenCV OCL module includes utility functions, low-level vision primitives, and high-level algorithms. The utility functions and low-level primitives provide a powerful infrastructure for developing fast vision algorithms taking advangtage of OCL whereas the high-level functionality (samples)includes some state-of-the-art algorithms (including LK Optical flow, and Face detection) ready to be used by the application developers. The module is also accompanied by an extensive performance and accuracy test suite.
The OpenCV OCL module includes utility functions, low-level vision primitives, and high-level algorithms. The utility functions and low-level primitives provide a powerful infrastructure for developing fast vision algorithms taking advantage of OCL, whereas the high-level functionality (samples) includes some state-of-the-art algorithms (including LK Optical flow, and Face detection) ready to be used by the application developers. The module is also accompanied by an extensive performance and accuracy test suite.
The OpenCV OCL module is designed for ease of use and does not require any knowledge of OpenCL. At a minimuml level, it can be viewed as a set of accelerators, that can take advantage of the high compute throughput that GPU/APU devices can provide. However, it can also be viewed as a starting point to really integratethe built-in functionality with your own custom OpenCL kernels, with or without modifying the source of OpenCV-OCL. Of course, knowledge of OpenCL will certainly help, however we hope that OpenCV-OCL module, and the kernels it contains in source code, can be very useful as a means of actually learning openCL. Such a knowledge would be necessary to further fine-tune any of the existing OpenCL kernels, or for extending the framework with new kernels. As of OpenCV 2.4.4, we introduce interoperability with OpenCL, enabling easy use of custom OpenCL kernels within the OpenCV framework.
The OpenCV OCL module is designed for ease of use and does not require any knowledge of OpenCL. At a minimum level, it can be viewed as a set of accelerators, that can take advantage of the high compute throughput that GPU/APU devices can provide. However, it can also be viewed as a starting point to really integrate the built-in functionality with your own custom OpenCL kernels, with or without modifying the source of OpenCV-OCL. Of course, knowledge of OpenCL will certainly help, however we hope that OpenCV-OCL module, and the kernels it contains in source code, can be very useful as a means of actually learning openCL. Such a knowledge would be necessary to further fine-tune any of the existing OpenCL kernels, or for extending the framework with new kernels. As of OpenCV 2.4.4, we introduce interoperability with OpenCL, enabling easy use of custom OpenCL kernels within the OpenCV framework.
To use the OCL module, you need to make sure that you have the OpenCL SDK provided with your device vendor. To correctly run the OCL module, you need to have the OpenCL runtime provide by the device vendor, typically the device driver.
To correctly run the OCL module, you need to have the OpenCL runtime provided by the device vendor, typically the device driver.
To enable OCL support, configure OpenCV using CMake with WITH\_OPENCL=ON. When the flag is set and if OpenCL SDK is installed, the full-featured OpenCV OCL module is built. Otherwise, the module may be not built. If you have AMD'S FFT and BLAS library, you can select it with WITH\_OPENCLAMDFFT=ON, WITH\_OPENCLAMDBLAS=ON.
To enable OCL support, configure OpenCV using CMake with ``WITH_OPENCL=ON``. When the flag is set and if OpenCL SDK is installed, the full-featured OpenCV OCL module is built. Otherwise, the module may be not built. If you have AMD'S FFT and BLAS library, you can select it with ``WITH_OPENCLAMDFFT=ON``, ``WITH_OPENCLAMDBLAS=ON``.
The ocl module can be found under the "modules" directory. In "modules/ocl/src" you can find the source code for the cpp class that wrap around the direct kernel invocation. The kernels themselves can be found in "modules/ocl/src/kernels." Samples can be found under "samples/ocl." Accuracy tests can be found in "modules/ocl/test," and performance tests under "module/ocl/perf."
The ocl module can be found under the "modules" directory. In "modules/ocl/src" you can find the source code for the cpp class that wrap around the direct kernel invocation. The kernels themselves can be found in "modules/ocl/src/opencl". Samples can be found under "samples/ocl". Accuracy tests can be found in "modules/ocl/test", and performance tests under "module/ocl/perf".
Right now, the user can select OpenCL device by specifying the environment variable ``OPENCV_OPENCL_DEVICE``. Variable format:
Right now, the user should define the cv::ocl::Info class in the application and call cv::ocl::getDevice before any cv::ocl::func. This operation initialize OpenCL runtime and set the first found device as computing device. If there are more than one device and you want to use undefault device, you can call cv::ocl::setDevice then.
.. code-block:: cpp
In the current version, all the thread share the same context and device so the multi-devices are not supported. We will add this feature soon. If a function support 4-channel operator, it should support 3-channel operator as well, because All the 3-channel matrix(i.e. RGB image) are represented by 4-channel matrix in oclMat. It means 3-channel image have 4-channel space with the last channel unused. We provide a transparent interface to handle the difference between OpenCV Mat and oclMat.
<Platform>:<CPU|GPU|ACCELERATOR|nothing=GPU/CPU>:<DeviceName or ID>
**Note:** Device ID range is: 0..9 (only one digit, 10 - it is a part of name)
Samples:
.. code-block:: cpp
'' = ':' = '::' = ':GPU|CPU:'
'AMD:GPU|CPU:'
'AMD::Tahiti'
':GPU:1'
':CPU:2'
Also the user can use ``cv::ocl::setDevice`` function (with ``cv::ocl::getOpenCLPlatforms`` and ``cv::ocl::getOpenCLDevices``). This function initializes OpenCL runtime and setup the passed device as computing device.
In the current version, all the thread share the same context and device so the multi-devices are not supported. We will add this feature soon. If a function support 4-channel operator, it should support 3-channel operator as well, because All the 3-channel matrix(i.e. RGB image) are represented by 4-channel matrix in ``oclMat``. It means 3-channel image have 4-channel space with the last channel unused. We provide a transparent interface to handle the difference between OpenCV Mat and ``oclMat``.
Developer Notes
-------------------
In a heterogeneous device environment, there may be cost associated with data transfer. This would be the case, for example, when data needs to be moved from host memory (accessible to the CPU), to device memory (accessible to a discrete GPU). in the case of integrated graphics chips, there may be performance issues, relating to memory coherency between access from the GPU "part" of the integrated device, or the CPU "part." For best performance, in either case, it is recommended that you do not introduce dat transfers between CPU and the discrete GPU, except in the beginning and the end of the algorithmic pipeline.
In a heterogeneous device environment, there may be cost associated with data transfer. This would be the case, for example, when data needs to be moved from host memory (accessible to the CPU), to device memory (accessible to a discrete GPU). in the case of integrated graphics chips, there may be performance issues, relating to memory coherency between access from the GPU "part" of the integrated device, or the CPU "part." For best performance, in either case, it is recommended that you do not introduce data transfers between CPU and the discrete GPU, except in the beginning and the end of the algorithmic pipeline.
Some tidbits:
1. OpenCL version should be larger than 1.1 with FULL PROFILE.
2. Currently (2.4.4) the user call the cv::ocl::getDevice before any other function in the ocl module. This will initialize the OpenCL runtime and set the first found device as computing device. If there are more than one device and you want to use undefault device, you can call cv::ocl::setDevice thereafter.
2. Currently there's only one OpenCL context and command queue. We hope to implement multi device and multi queue support in the future.
3. Many kernels use 256 as its workgroup size if possible, so the max work group size of the device must larger than 256. All GPU devices we are aware of indeed support 256 workitems in a workgroup, however non GPU devices may not. This will be improved in the future.
4. If the device does not support double arithetic, we revert to float.
4. If the device does not support double arithmetic, then functions' implementation generates an error.
5. The oclMat uses buffer object, not image object.
5. The ``oclMat`` uses buffer object, not image object.
6. All the 3-channel matrices(i.e. RGB image) are represented by 4-channel matrices in oclMat, with the last channel unused. We provide a transparent interface to handle the difference between OpenCV Mat and oclMat.
6. All the 3-channel matrices (i.e. RGB image) are represented by 4-channel matrices in ``oclMat``, with the last channel unused. We provide a transparent interface to handle the difference between OpenCV Mat and ``oclMat``.
7. All the matrix in oclMat is aligned in column(now the alignment factor is 32 byte). It means, if a matrix is n columns m rows with the element size x byte, we will assign ALIGNMENT(x*n) bytes for each column with the last ALIGNMENT(x*n) - x*n bytes unused, so there's small holes after each column if its size is not the multiply of ALIGN.
7. All the matrix in ``oclMat`` is aligned in column (now the alignment factor for ``step`` is 32+ byte). It means, m.cols * m.elemSize() <= m.step.
8. Data transfer between Mat and oclMat: If the CPU matrix is aligned in column, we will use faster API to transfer between Mat and oclMat, otherwise, we will use clEnqueueRead/WriteBufferRect to transfer data to guarantee the alignment. 3-channel matrix is an exception, it's directly transferred to a temp buffer and then padded to 4-channel matrix(also aligned) when uploading and do the reverse operation when downloading.
8. Data transfer between Mat and ``oclMat``: If the CPU matrix is aligned in column, we will use faster API to transfer between Mat and ``oclMat``, otherwise, we will use clEnqueueRead/WriteBufferRect to transfer data to guarantee the alignment. 3-channel matrix is an exception, it's directly transferred to a temp buffer and then padded to 4-channel matrix(also aligned) when uploading and do the reverse operation when downloading.
9. Data transfer between Mat and oclMat: ROI is a feature of OpenCV, which allow users process a sub rectangle of a matrix. When a CPU matrix which has ROI will be transfered to GPU, the whole matrix will be transfered and set ROI as CPU's. In a word, we always transfer the whole matrix despite whether it has ROI or not.
9. Data transfer between Mat and ``oclMat``: ROI is a feature of OpenCV, which allow users process a sub rectangle of a matrix. When a CPU matrix which has ROI will be transfered to GPU, the whole matrix will be transfered and set ROI as CPU's. In a word, we always transfer the whole matrix despite whether it has ROI or not.
10. All the kernel file should locate in ocl/src/kernels/ with the extension ".cl". ALL the kernel files are transformed to pure characters at compilation time in kernels.cpp, and the file name without extension is the name of the characters.
10. All the kernel file should locate in "modules/ocl/src/opencl/" with the extension ".cl". All the kernel files are transformed to pure characters at compilation time in opencl_kernels.cpp, and the file name without extension is the name of the program sources.

@ -91,11 +91,7 @@ ocl::distanceToCenters
----------------------
For each samples in ``source``, find its closest neighour in ``centers``.
.. ocv:function:: void ocl::distanceToCenters(oclMat &dists, oclMat &labels, const oclMat &src, const oclMat &centers, int distType = NORM_L2SQR, const oclMat &indices = oclMat())
:param dists: The output distances calculated from each sample to the best matched center.
:param labels: The output index of best matched center for each row of sample.
.. ocv:function:: void ocl::distanceToCenters(const oclMat &src, const oclMat &centers, Mat &dists, Mat &labels, int distType = NORM_L2SQR)
:param src: Floating-point matrix of input samples. One row per sample.
@ -103,10 +99,8 @@ For each samples in ``source``, find its closest neighour in ``centers``.
:param distType: Distance metric to calculate distances. Supports ``NORM_L1`` and ``NORM_L2SQR``.
:param indices: Optional source indices. If not empty:
:param dists: The output distances calculated from each sample to the best matched center.
* only the indexed source samples will be processed
* outputs, i.e., ``dists`` and ``labels``, have the same size of indices
* outputs are in the same order of indices instead of the order of src
:param labels: The output index of best matched center for each row of sample.
The method is a utility function which maybe used for multiple clustering algorithms such as K-means.

@ -117,7 +117,6 @@ Computes a dense optical flow using the Gunnar Farneback's algorithm.
:param frame1: Second 8-bit gray-scale input image
:param flowx: Flow horizontal component
:param flowy: Flow vertical component
:param s: Stream
.. seealso:: :ocv:func:`calcOpticalFlowFarneback`
@ -230,8 +229,6 @@ Interpolates frames (images) using provided optical flow (displacement field).
:param buf: Temporary buffer, will have width x 6*height size, CV_32FC1 type and contain 6 oclMat: occlusion masks for first frame, occlusion masks for second, interpolated forward horizontal flow, interpolated forward vertical flow, interpolated backward horizontal flow, interpolated backward vertical flow.
:param stream: Stream for the asynchronous version.
ocl::KalmanFilter
--------------------
.. ocv:class:: ocl::KalmanFilter
@ -418,8 +415,6 @@ Updates the background model and returns the foreground mask.
:param fgmask: The output foreground mask as an 8-bit binary image.
:param stream: Stream for the asynchronous version.
ocl::MOG::getBackgroundImage
--------------------------------
@ -429,8 +424,6 @@ Computes a background image.
:param backgroundImage: The output background image.
:param stream: Stream for the asynchronous version.
ocl::MOG::release
---------------------
@ -443,7 +436,9 @@ ocl::MOG2
-------------
.. ocv:class:: ocl::MOG2 : public ocl::BackgroundSubtractor
Gaussian Mixture-based Background/Foreground Segmentation Algorithm. ::
Gaussian Mixture-based Background/Foreground Segmentation Algorithm.
The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [MOG2004]_. ::
class CV_EXPORTS MOG2: public cv::ocl::BackgroundSubtractor
{
@ -485,45 +480,42 @@ Gaussian Mixture-based Background/Foreground Segmentation Algorithm. ::
/* hidden */
};
The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [MOG2004]_.
.. ocv:member:: float backgroundRatio
Here are important members of the class that control the algorithm, which you can set after constructing the class instance:
Threshold defining whether the component is significant enough to be included into the background model. ``cf=0.1 => TB=0.9`` is default. For ``alpha=0.001``, it means that the mode should exist for approximately 105 frames before it is considered foreground.
.. ocv:member:: float backgroundRatio
.. ocv:member:: float varThreshold
Threshold defining whether the component is significant enough to be included into the background model. ``cf=0.1 => TB=0.9`` is default. For ``alpha=0.001``, it means that the mode should exist for approximately 105 frames before it is considered foreground.
Threshold for the squared Mahalanobis distance that helps decide when a sample is close to the existing components (corresponds to ``Tg``). If it is not close to any component, a new component is generated. ``3 sigma => Tg=3*3=9`` is default. A smaller ``Tg`` value generates more components. A higher ``Tg`` value may result in a small number of components but they can grow too large.
.. ocv:member:: float varThreshold
.. ocv:member:: float fVarInit
Threshold for the squared Mahalanobis distance that helps decide when a sample is close to the existing components (corresponds to ``Tg``). If it is not close to any component, a new component is generated. ``3 sigma => Tg=3*3=9`` is default. A smaller ``Tg`` value generates more components. A higher ``Tg`` value may result in a small number of components but they can grow too large.
Initial variance for the newly generated components. It affects the speed of adaptation. The parameter value is based on your estimate of the typical standard deviation from the images. OpenCV uses 15 as a reasonable value.
.. ocv:member:: float fVarInit
.. ocv:member:: float fVarMin
Initial variance for the newly generated components. It affects the speed of adaptation. The parameter value is based on your estimate of the typical standard deviation from the images. OpenCV uses 15 as a reasonable value.
Parameter used to further control the variance.
.. ocv:member:: float fVarMin
.. ocv:member:: float fVarMax
Parameter used to further control the variance.
Parameter used to further control the variance.
.. ocv:member:: float fVarMax
.. ocv:member:: float fCT
Parameter used to further control the variance.
Complexity reduction parameter. This parameter defines the number of samples needed to accept to prove the component exists. ``CT=0.05`` is a default value for all the samples. By setting ``CT=0`` you get an algorithm very similar to the standard Stauffer&Grimson algorithm.
.. ocv:member:: float fCT
.. ocv:member:: uchar nShadowDetection
Complexity reduction parameter. This parameter defines the number of samples needed to accept to prove the component exists. ``CT=0.05`` is a default value for all the samples. By setting ``CT=0`` you get an algorithm very similar to the standard Stauffer&Grimson algorithm.
The value for marking shadow pixels in the output foreground mask. Default value is 127.
.. ocv:member:: uchar nShadowDetection
.. ocv:member:: float fTau
The value for marking shadow pixels in the output foreground mask. Default value is 127.
Shadow threshold. The shadow is detected if the pixel is a darker version of the background. ``Tau`` is a threshold defining how much darker the shadow can be. ``Tau= 0.5`` means that if a pixel is more than twice darker then it is not shadow. See [ShadowDetect2003]_.
.. ocv:member:: float fTau
.. ocv:member:: bool bShadowDetection
Shadow threshold. The shadow is detected if the pixel is a darker version of the background. ``Tau`` is a threshold defining how much darker the shadow can be. ``Tau= 0.5`` means that if a pixel is more than twice darker then it is not shadow. See [ShadowDetect2003]_.
Parameter defining whether shadow detection should be enabled.
.. ocv:member:: bool bShadowDetection
Parameter defining whether shadow detection should be enabled.
.. seealso:: :ocv:class:`BackgroundSubtractorMOG2`
@ -549,8 +541,6 @@ Updates the background model and returns the foreground mask.
:param fgmask: The output foreground mask as an 8-bit binary image.
:param stream: Stream for the asynchronous version.
ocl::MOG2::getBackgroundImage
---------------------------------
@ -560,11 +550,9 @@ Computes a background image.
:param backgroundImage: The output background image.
:param stream: Stream for the asynchronous version.
ocl::MOG2::release
----------------------
Releases all inner buffer's memory.
.. ocv:function:: void ocl::MOG2::release()
.. ocv:function:: void ocl::MOG2::release()

@ -305,16 +305,13 @@ namespace cv
void copyTo( oclMat &m, const oclMat &mask = oclMat()) const;
//! converts oclMatrix to another datatype with optional scalng. See cvConvertScale.
//It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
void convertTo( oclMat &m, int rtype, double alpha = 1, double beta = 0 ) const;
void assignTo( oclMat &m, int type = -1 ) const;
//! sets every oclMatrix element to s
//It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
oclMat& operator = (const Scalar &s);
//! sets some of the oclMatrix elements to s, according to the mask
//It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
oclMat& setTo(const Scalar &s, const oclMat &mask = oclMat());
//! creates alternative oclMatrix header for the same data, with different
// number of channels and/or different number of rows. see cvReshape.
@ -555,11 +552,12 @@ namespace cv
CV_EXPORTS void bilateralFilter(const oclMat& src, oclMat& dst, int d, double sigmaColor, double sigmaSpace, int borderType=BORDER_DEFAULT);
//! Applies an adaptive bilateral filter to the input image
// This is not truly a bilateral filter. Instead of using user provided fixed parameters,
// the function calculates a constant at each window based on local standard deviation,
// and use this constant to do filtering.
// Unlike the usual bilateral filter that uses fixed value for sigmaColor,
// the adaptive version calculates the local variance in he ksize neighborhood
// and use this as sigmaColor, for the value filtering. However, the local standard deviation is
// clamped to the maxSigmaColor.
// supports 8UC1, 8UC3
CV_EXPORTS void adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, Point anchor = Point(-1, -1), int borderType=BORDER_DEFAULT);
CV_EXPORTS void adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, double maxSigmaColor=20.0, Point anchor = Point(-1, -1), int borderType=BORDER_DEFAULT);
//! computes exponent of each matrix element (dst = e**src)
// supports only CV_32FC1, CV_64FC1 type
@ -650,6 +648,9 @@ namespace cv
//! initializes a scaled identity matrix
CV_EXPORTS void setIdentity(oclMat& src, const Scalar & val = Scalar(1));
//! fills the output array with repeated copies of the input array
CV_EXPORTS void repeat(const oclMat & src, int ny, int nx, oclMat & dst);
//////////////////////////////// Filter Engine ////////////////////////////////
/*!
@ -898,7 +899,7 @@ namespace cv
// supports NORM_L1 and NORM_L2 distType
// if indices is provided, only the indexed rows will be calculated and their results are in the same
// order of indices
CV_EXPORTS void distanceToCenters(oclMat &dists, oclMat &labels, const oclMat &src, const oclMat &centers, int distType = NORM_L2SQR, const oclMat &indices = oclMat());
CV_EXPORTS void distanceToCenters(const oclMat &src, const oclMat &centers, Mat &dists, Mat &labels, int distType = NORM_L2SQR);
//!Does k-means procedure on GPU
// supports CV_32FC1/CV_32FC2/CV_32FC4 data type
@ -913,8 +914,8 @@ namespace cv
{
public:
void detectMultiScale(oclMat &image, CV_OUT std::vector<cv::Rect>& faces,
double scaleFactor = 1.1, int minNeighbors = 3, int flags = 0,
Size minSize = Size(), Size maxSize = Size());
double scaleFactor = 1.1, int minNeighbors = 3, int flags = 0,
Size minSize = Size(), Size maxSize = Size());
};
/////////////////////////////// Pyramid /////////////////////////////////////

@ -39,16 +39,16 @@
//
//M*/
#if !defined(DUMP_INFO_STDOUT) && !defined(DUMP_INFO_XML)
#if !defined(DUMP_MESSAGE_STDOUT) && !defined(DUMP_PROPERTY_XML)
#error Invalid usage
#endif
#if !defined(DUMP_INFO_STDOUT)
#define DUMP_INFO_STDOUT(...)
#if !defined(DUMP_PROPERTY_XML)
#define DUMP_PROPERTY_XML(...)
#endif
#if !defined(DUMP_INFO_XML)
#define DUMP_INFO_XML(...)
#if !defined(DUMP_MESSAGE_STDOUT)
#define DUMP_MESSAGE_STDOUT(...)
#endif
#include <sstream>
@ -85,51 +85,70 @@ static void dumpOpenCLDevice()
using namespace cv::ocl;
try
{
cv::ocl::PlatformsInfo platforms;
cv::ocl::getOpenCLPlatforms(platforms);
DUMP_MESSAGE_STDOUT("OpenCL Platforms: ");
const char* deviceTypeStr;
for(unsigned int i=0; i < platforms.size(); i++)
{
DUMP_MESSAGE_STDOUT(" " << platforms.at(i)->platformName);
const cv::ocl::DevicesInfo& devices = platforms.at(i)->devices;
for(unsigned int j=0; j < devices.size(); j++)
{
const cv::ocl::DeviceInfo& current_device = *devices.at(j);
deviceTypeStr = current_device.deviceType == CVCL_DEVICE_TYPE_CPU
? ("CPU") : (current_device.deviceType == CVCL_DEVICE_TYPE_GPU ? "GPU" : "unknown");
DUMP_MESSAGE_STDOUT( " " << deviceTypeStr << " : " << current_device.deviceName << " : " << current_device.deviceVersion );
DUMP_PROPERTY_XML("cv_ocl_platform_"<< i<<"_device_"<<j, "(Platform=" << current_device.platform->platformName << ")(Type="
<< deviceTypeStr <<")(Name="<< current_device.deviceName << ")(Version="<< current_device.deviceVersion<<")");
}
}
DUMP_MESSAGE_STDOUT("Current OpenCL device: ");
const cv::ocl::DeviceInfo& deviceInfo = cv::ocl::Context::getContext()->getDeviceInfo();
const char* deviceTypeStr = deviceInfo.deviceType == CVCL_DEVICE_TYPE_CPU
? "CPU" :
(deviceInfo.deviceType == CVCL_DEVICE_TYPE_GPU ? "GPU" : "unknown");
DUMP_INFO_STDOUT("Device type", deviceTypeStr);
DUMP_INFO_XML("cv_ocl_deviceType", deviceTypeStr);
DUMP_MESSAGE_STDOUT(" Platform = "<< deviceInfo.platform->platformName);
DUMP_PROPERTY_XML("cv_ocl_current_platformName", deviceInfo.platform->platformName);
DUMP_INFO_STDOUT("Platform name", deviceInfo.platform->platformName);
DUMP_INFO_XML("cv_ocl_platformName", deviceInfo.platform->platformName);
deviceTypeStr = deviceInfo.deviceType == CVCL_DEVICE_TYPE_CPU
? "CPU" : (deviceInfo.deviceType == CVCL_DEVICE_TYPE_GPU ? "GPU" : "unknown");
DUMP_MESSAGE_STDOUT(" Type = "<< deviceTypeStr);
DUMP_PROPERTY_XML("cv_ocl_current_deviceType", deviceTypeStr);
DUMP_INFO_STDOUT("Device name", deviceInfo.deviceName);
DUMP_INFO_XML("cv_ocl_deviceName", deviceInfo.deviceName);
DUMP_MESSAGE_STDOUT(" Name = "<< deviceInfo.deviceName);
DUMP_PROPERTY_XML("cv_ocl_current_deviceName", deviceInfo.deviceName);
DUMP_INFO_STDOUT("Device version", deviceInfo.deviceVersion);
DUMP_INFO_XML("cv_ocl_deviceVersion", deviceInfo.deviceVersion);
DUMP_MESSAGE_STDOUT(" Version = " << deviceInfo.deviceVersion);
DUMP_PROPERTY_XML("cv_ocl_current_deviceVersion", deviceInfo.deviceVersion);
DUMP_INFO_STDOUT("Compute units", deviceInfo.maxComputeUnits);
DUMP_INFO_XML("cv_ocl_maxComputeUnits", deviceInfo.maxComputeUnits);
DUMP_MESSAGE_STDOUT(" Compute units = "<< deviceInfo.maxComputeUnits);
DUMP_PROPERTY_XML("cv_ocl_current_maxComputeUnits", deviceInfo.maxComputeUnits);
DUMP_INFO_STDOUT("Max work group size", deviceInfo.maxWorkGroupSize);
DUMP_INFO_XML("cv_ocl_maxWorkGroupSize", deviceInfo.maxWorkGroupSize);
DUMP_MESSAGE_STDOUT(" Max work group size = "<< deviceInfo.maxWorkGroupSize);
DUMP_PROPERTY_XML("cv_ocl_current_maxWorkGroupSize", deviceInfo.maxWorkGroupSize);
std::string localMemorySizeStr = bytesToStringRepr(deviceInfo.localMemorySize);
DUMP_INFO_STDOUT("Local memory size", localMemorySizeStr.c_str());
DUMP_INFO_XML("cv_ocl_localMemorySize", deviceInfo.localMemorySize);
DUMP_MESSAGE_STDOUT(" Local memory size = "<< localMemorySizeStr.c_str());
DUMP_PROPERTY_XML("cv_ocl_current_localMemorySize", deviceInfo.localMemorySize);
std::string maxMemAllocSizeStr = bytesToStringRepr(deviceInfo.maxMemAllocSize);
DUMP_INFO_STDOUT("Max memory allocation size", maxMemAllocSizeStr.c_str());
DUMP_INFO_XML("cv_ocl_maxMemAllocSize", deviceInfo.maxMemAllocSize);
DUMP_MESSAGE_STDOUT(" Max memory allocation size = "<< maxMemAllocSizeStr.c_str());
DUMP_PROPERTY_XML("cv_ocl_current_maxMemAllocSize", deviceInfo.maxMemAllocSize);
const char* doubleSupportStr = deviceInfo.haveDoubleSupport ? "Yes" : "No";
DUMP_INFO_STDOUT("Double support", doubleSupportStr);
DUMP_INFO_XML("cv_ocl_haveDoubleSupport", deviceInfo.haveDoubleSupport);
DUMP_MESSAGE_STDOUT(" Double support = "<< doubleSupportStr);
DUMP_PROPERTY_XML("cv_ocl_current_haveDoubleSupport", deviceInfo.haveDoubleSupport);
const char* isUnifiedMemoryStr = deviceInfo.isUnifiedMemory ? "Yes" : "No";
DUMP_INFO_STDOUT("Unified memory", isUnifiedMemoryStr);
DUMP_INFO_XML("cv_ocl_isUnifiedMemory", deviceInfo.isUnifiedMemory);
DUMP_MESSAGE_STDOUT(" Unified memory = "<< isUnifiedMemoryStr);
DUMP_PROPERTY_XML("cv_ocl_current_isUnifiedMemory", deviceInfo.isUnifiedMemory);
}
catch (...)
{
DUMP_INFO_STDOUT("OpenCL device", "not available");
DUMP_INFO_XML("cv_ocl", "not available");
DUMP_MESSAGE_STDOUT("OpenCL device not available");
DUMP_PROPERTY_XML("cv_ocl", "not available");
}
}
#undef DUMP_INFO_STDOUT
#undef DUMP_INFO_XML
#undef DUMP_MESSAGE_STDOUT
#undef DUMP_PROPERTY_XML

@ -42,17 +42,20 @@
#include "perf_precomp.hpp"
#define DUMP_INFO_STDOUT(propertyDisplayName, propertyValue) \
#define DUMP_PROPERTY_XML(propertyName, propertyValue) \
do { \
std::cout << (propertyDisplayName) << ": " << (propertyValue) << std::endl; \
std::stringstream ssName, ssValue;\
ssName << propertyName;\
ssValue << propertyValue; \
::testing::Test::RecordProperty(ssName.str(), ssValue.str()); \
} while (false)
#define DUMP_INFO_XML(propertyXMLName, propertyValue) \
#define DUMP_MESSAGE_STDOUT(msg) \
do { \
std::stringstream ss; ss << propertyValue; \
::testing::Test::RecordProperty((propertyXMLName), ss.str()); \
std::cout << msg << std::endl; \
} while (false)
#include "opencv2/ocl/private/opencl_dumpinfo.hpp"
static const char * impls[] =

@ -1051,3 +1051,40 @@ PERF_TEST_P(AbsFixture, Abs,
else
OCL_PERF_ELSE
}
///////////// Repeat ////////////////////////
typedef Size_MatType RepeatFixture;
PERF_TEST_P(RepeatFixture, Repeat,
::testing::Combine(::testing::Values(OCL_SIZE_1000, OCL_SIZE_2000),
OCL_PERF_ENUM(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4)))
{
const Size_MatType_t params = GetParam();
const Size srcSize = get<0>(params);
const int type = get<1>(params);
const int nx = 3, ny = 2;
const Size dstSize(srcSize.width * nx, srcSize.height * ny);
Mat src(srcSize, type), dst(dstSize, type);
declare.in(src, WARMUP_RNG).out(dst);
if (RUN_OCL_IMPL)
{
ocl::oclMat oclSrc(src), oclDst(dstSize, type);
OCL_TEST_CYCLE() cv::ocl::repeat(oclSrc, ny, nx, oclDst);
oclDst.download(dst);
SANITY_CHECK(dst);
}
else if (RUN_PLAIN_IMPL)
{
TEST_CYCLE() cv::repeat(src, ny, nx, dst);
SANITY_CHECK(dst);
}
else
OCL_PERF_ELSE
}

@ -50,7 +50,6 @@
using namespace perf;
///////////// Haar ////////////////////////
PERF_TEST(HaarFixture, Haar)
{
vector<Rect> faces;

@ -872,58 +872,57 @@ PERF_TEST_P(columnSumFixture, columnSum, OCL_TYPICAL_MAT_SIZES)
//////////////////////////////distanceToCenters////////////////////////////////////////////////
CV_ENUM(DistType, NORM_L1, NORM_L2SQR);
CV_ENUM(DistType, NORM_L1, NORM_L2SQR)
typedef tuple<Size, DistType> distanceToCentersParameters;
typedef TestBaseWithParam<distanceToCentersParameters> distanceToCentersFixture;
static void distanceToCentersPerfTest(Mat& src, Mat& centers, Mat& dists, Mat& labels, int distType)
{
Mat batch_dists;
cv::batchDistance(src,centers,batch_dists, CV_32FC1, noArray(), distType);
cv::batchDistance(src, centers, batch_dists, CV_32FC1, noArray(), distType);
std::vector<float> dists_v;
std::vector<int> labels_v;
for(int i = 0; i<batch_dists.rows; i++)
for (int i = 0; i < batch_dists.rows; i++)
{
Mat r = batch_dists.row(i);
double mVal;
Point mLoc;
minMaxLoc(r, &mVal, NULL, &mLoc, NULL);
dists_v.push_back((float)mVal);
dists_v.push_back(static_cast<float>(mVal));
labels_v.push_back(mLoc.x);
}
Mat temp_dists(dists_v);
Mat temp_labels(labels_v);
temp_dists.reshape(1,1).copyTo(dists);
temp_labels.reshape(1,1).copyTo(labels);
Mat(dists_v).copyTo(dists);
Mat(labels_v).copyTo(labels);
}
PERF_TEST_P(distanceToCentersFixture, distanceToCenters, ::testing::Combine(::testing::Values(cv::Size(256,256), cv::Size(512,512)), DistType::all()) )
{
Size size = get<0>(GetParam());
int distType = get<1>(GetParam());
Mat src(size, CV_32FC1);
Mat centers(size, CV_32FC1);
Mat dists(cv::Size(src.rows,1), CV_32FC1);
Mat labels(cv::Size(src.rows,1), CV_32SC1);
Mat src(size, CV_32FC1), centers(size, CV_32FC1);
Mat dists(src.rows, 1, CV_32FC1), labels(src.rows, 1, CV_32SC1);
declare.in(src, centers, WARMUP_RNG).out(dists, labels);
if (RUN_OCL_IMPL)
{
ocl::oclMat ocl_src(src);
ocl::oclMat ocl_centers(centers);
ocl::oclMat ocl_dists(dists);
ocl::oclMat ocl_labels(labels);
ocl::oclMat ocl_src(src), ocl_centers(centers);
OCL_TEST_CYCLE() ocl::distanceToCenters(ocl_dists,ocl_labels,ocl_src, ocl_centers, distType);
ocl_dists.download(dists);
ocl_labels.download(labels);
OCL_TEST_CYCLE() ocl::distanceToCenters(ocl_src, ocl_centers, dists, labels, distType);
SANITY_CHECK(dists, 1e-6, ERROR_RELATIVE);
SANITY_CHECK(labels);
}
else if (RUN_PLAIN_IMPL)
{
TEST_CYCLE() distanceToCentersPerfTest(src,centers,dists,labels,distType);
TEST_CYCLE() distanceToCentersPerfTest(src, centers, dists, labels, distType);
SANITY_CHECK(dists, 1e-6, ERROR_RELATIVE);
SANITY_CHECK(labels);
}

@ -1706,3 +1706,21 @@ void cv::ocl::setIdentity(oclMat& src, const Scalar & scalar)
openCLExecuteKernel(src.clCxt, &arithm_setidentity, "setIdentity", global_threads, local_threads,
args, -1, -1, buildOptions.c_str());
}
//////////////////////////////////////////////////////////////////////////////
////////////////////////////////// Repeat ////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void cv::ocl::repeat(const oclMat & src, int ny, int nx, oclMat & dst)
{
CV_Assert(nx > 0 && ny > 0);
dst.create(src.rows * ny, src.cols * nx, src.type());
for (int y = 0; y < ny; ++y)
for (int x = 0; x < nx; ++x)
{
Rect roi(x * src.cols, y * src.rows, src.cols, src.rows);
oclMat hdr = dst(roi);
src.copyTo(hdr);
}
}

@ -109,12 +109,15 @@ cl_mem openCLCreateBuffer(Context *ctx, size_t flag , size_t size)
return buffer;
}
#define MEMORY_CORRUPTION_GUARD
#ifdef MEMORY_CORRUPTION_GUARD
//#define CHECK_MEMORY_CORRUPTION
#ifdef CHECK_MEMORY_CORRUPTION
//#define CHECK_MEMORY_CORRUPTION_PRINT_ERROR
#define CHECK_MEMORY_CORRUPTION_PRINT_ERROR
#define CHECK_MEMORY_CORRUPTION_RAISE_ERROR
static const int __memory_corruption_check_bytes = 1024*1024;
static const int __memory_corruption_guard_bytes = 64*1024;
#ifdef CHECK_MEMORY_CORRUPTION
static const int __memory_corruption_check_pattern = 0x14326547; // change pattern for sizeof(int)==8
#endif
struct CheckBuffers
{
cl_mem mainBuffer;
@ -128,7 +131,7 @@ struct CheckBuffers
CheckBuffers(cl_mem _mainBuffer, size_t _size, size_t _widthInBytes, size_t _height)
: mainBuffer(_mainBuffer), size(_size), widthInBytes(_widthInBytes), height(_height)
{
// notihng
// nothing
}
};
static std::map<cl_mem, CheckBuffers> __check_buffers;
@ -145,32 +148,52 @@ void openCLMallocPitchEx(Context *ctx, void **dev_ptr, size_t *pitch,
{
cl_int status;
size_t size = widthInBytes * height;
#ifndef CHECK_MEMORY_CORRUPTION
*dev_ptr = clCreateBuffer(getClContext(ctx), gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
size, 0, &status);
openCLVerifyCall(status);
bool useSubBuffers =
#ifndef MEMORY_CORRUPTION_GUARD
false;
#else
size_t allocSize = size + __memory_corruption_check_bytes * 2;
cl_mem mainBuffer = clCreateBuffer(getClContext(ctx), gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
allocSize, 0, &status);
openCLVerifyCall(status);
cl_buffer_region r = {__memory_corruption_check_bytes, size};
*dev_ptr = clCreateSubBuffer(mainBuffer,
gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
CL_BUFFER_CREATE_TYPE_REGION, &r,
&status);
openCLVerifyCall(status);
std::vector<int> tmp(__memory_corruption_check_bytes / sizeof(int),
__memory_corruption_check_pattern);
CV_Assert(tmp.size() * sizeof(int) == __memory_corruption_check_bytes);
openCLVerifyCall(clEnqueueWriteBuffer(getClCommandQueue(ctx),
mainBuffer, CL_TRUE, 0, __memory_corruption_check_bytes, &tmp[0],
0, NULL, NULL));
openCLVerifyCall(clEnqueueWriteBuffer(getClCommandQueue(ctx),
mainBuffer, CL_TRUE, __memory_corruption_check_bytes + size, __memory_corruption_check_bytes, &tmp[0],
0, NULL, NULL));
CheckBuffers data(mainBuffer, size, widthInBytes, height);
__check_buffers.insert(std::pair<cl_mem, CheckBuffers>((cl_mem)*dev_ptr, data));
true;
#endif
const DeviceInfo& devInfo = ctx->getDeviceInfo();
if (useSubBuffers && devInfo.isIntelDevice)
{
useSubBuffers = false; // TODO FIXIT We observe memory leaks then we working with sub-buffers
// on the CPU device of Intel OpenCL SDK (Linux). We will investigate this later.
}
if (!useSubBuffers)
{
*dev_ptr = clCreateBuffer(getClContext(ctx), gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
size, 0, &status);
openCLVerifyCall(status);
}
#ifdef MEMORY_CORRUPTION_GUARD
else
{
size_t allocSize = size + __memory_corruption_guard_bytes * 2;
cl_mem mainBuffer = clCreateBuffer(getClContext(ctx), gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
allocSize, 0, &status);
openCLVerifyCall(status);
cl_buffer_region r = {__memory_corruption_guard_bytes, size};
*dev_ptr = clCreateSubBuffer(mainBuffer,
gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
CL_BUFFER_CREATE_TYPE_REGION, &r,
&status);
openCLVerifyCall(status);
#ifdef CHECK_MEMORY_CORRUPTION
std::vector<int> tmp(__memory_corruption_guard_bytes / sizeof(int),
__memory_corruption_check_pattern);
CV_Assert(tmp.size() * sizeof(int) == __memory_corruption_guard_bytes);
openCLVerifyCall(clEnqueueWriteBuffer(getClCommandQueue(ctx),
mainBuffer, CL_FALSE, 0, __memory_corruption_guard_bytes, &tmp[0],
0, NULL, NULL));
openCLVerifyCall(clEnqueueWriteBuffer(getClCommandQueue(ctx),
mainBuffer, CL_FALSE, __memory_corruption_guard_bytes + size, __memory_corruption_guard_bytes, &tmp[0],
0, NULL, NULL));
clFinish(getClCommandQueue(ctx));
#endif
CheckBuffers data(mainBuffer, size, widthInBytes, height);
__check_buffers.insert(std::pair<cl_mem, CheckBuffers>((cl_mem)*dev_ptr, data));
}
#endif
*pitch = widthInBytes;
}
@ -224,40 +247,48 @@ void openCLCopyBuffer2D(Context *ctx, void *dst, size_t dpitch, int dst_offset,
void openCLFree(void *devPtr)
{
openCLSafeCall(clReleaseMemObject((cl_mem)devPtr));
#ifdef MEMORY_CORRUPTION_GUARD
#ifdef CHECK_MEMORY_CORRUPTION
bool failBefore = false, failAfter = false;
#endif
CheckBuffers data;
std::map<cl_mem, CheckBuffers>::iterator i = __check_buffers.find((cl_mem)devPtr);
if (i != __check_buffers.end())
{
data = i->second;
#ifdef CHECK_MEMORY_CORRUPTION
Context* ctx = Context::getContext();
std::vector<uchar> checkBefore(__memory_corruption_check_bytes);
std::vector<uchar> checkAfter(__memory_corruption_check_bytes);
std::vector<uchar> checkBefore(__memory_corruption_guard_bytes);
std::vector<uchar> checkAfter(__memory_corruption_guard_bytes);
openCLVerifyCall(clEnqueueReadBuffer(getClCommandQueue(ctx),
data.mainBuffer, CL_TRUE, 0, __memory_corruption_check_bytes, &checkBefore[0],
data.mainBuffer, CL_FALSE, 0, __memory_corruption_guard_bytes, &checkBefore[0],
0, NULL, NULL));
openCLVerifyCall(clEnqueueReadBuffer(getClCommandQueue(ctx),
data.mainBuffer, CL_TRUE, __memory_corruption_check_bytes + data.size, __memory_corruption_check_bytes, &checkAfter[0],
data.mainBuffer, CL_FALSE, __memory_corruption_guard_bytes + data.size, __memory_corruption_guard_bytes, &checkAfter[0],
0, NULL, NULL));
clFinish(getClCommandQueue(ctx));
std::vector<int> tmp(__memory_corruption_check_bytes / sizeof(int),
std::vector<int> tmp(__memory_corruption_guard_bytes / sizeof(int),
__memory_corruption_check_pattern);
if (memcmp(&checkBefore[0], &tmp[0], __memory_corruption_check_bytes) != 0)
if (memcmp(&checkBefore[0], &tmp[0], __memory_corruption_guard_bytes) != 0)
{
failBefore = true;
}
if (memcmp(&checkAfter[0], &tmp[0], __memory_corruption_check_bytes) != 0)
if (memcmp(&checkAfter[0], &tmp[0], __memory_corruption_guard_bytes) != 0)
{
failAfter = true;
}
#else
// TODO FIXIT Attach clReleaseMemObject call to event completion callback
Context* ctx = Context::getContext();
clFinish(getClCommandQueue(ctx));
#endif
openCLSafeCall(clReleaseMemObject(data.mainBuffer));
__check_buffers.erase(i);
}
#endif
openCLSafeCall(clReleaseMemObject((cl_mem)devPtr));
#ifdef CHECK_MEMORY_CORRUPTION
#if defined(CHECK_MEMORY_CORRUPTION)
if (failBefore)
{
#ifdef CHECK_MEMORY_CORRUPTION_PRINT_ERROR
@ -276,7 +307,8 @@ void openCLFree(void *devPtr)
CV_Error(CV_StsInternal, "Memory corruption detected: after buffer");
#endif
}
#endif
#endif // CHECK_MEMORY_CORRUPTION
#endif // MEMORY_CORRUPTION_GUARD
}
cl_kernel openCLGetKernelFromSource(const Context *ctx, const cv::ocl::ProgramEntry* source, String kernelName)

@ -20,6 +20,7 @@
// Zero Lin, Zero.Lin@amd.com
// Zhang Ying, zhangying913@gmail.com
// Yao Wang, bitwangyaoyao@gmail.com
// Harris Gasparakis, harris.gasparakis@amd.com
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
@ -1409,7 +1410,7 @@ void cv::ocl::GaussianBlur(const oclMat &src, oclMat &dst, Size ksize, double si
////////////////////////////////////////////////////////////////////////////////////////////////////
// Adaptive Bilateral Filter
void cv::ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, Point anchor, int borderType)
void cv::ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, double maxSigmaColor, Point anchor, int borderType)
{
CV_Assert((ksize.width & 1) && (ksize.height & 1)); // ksize must be odd
CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC3); // source must be 8bit RGB image
@ -1420,10 +1421,24 @@ void cv::ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize
int idx = 0;
int w = ksize.width / 2;
int h = ksize.height / 2;
for(int y=-h; y<=h; y++)
for(int x=-w; x<=w; x++)
int ABF_GAUSSIAN_ocl = 1;
if(ABF_GAUSSIAN_ocl)
{
lut.at<float>(idx++) = sigma2 / (sigma2 + x * x + y * y);
for(int y=-h; y<=h; y++)
for(int x=-w; x<=w; x++)
{
lut.at<float>(idx++) = expf( (float)(-0.5 * (x * x + y * y)/sigma2));
}
}
else
{
for(int y=-h; y<=h; y++)
for(int x=-w; x<=w; x++)
{
lut.at<float>(idx++) = (float) (sigma2 / (sigma2 + x * x + y * y));
}
}
oclMat dlut(lut);
@ -1431,7 +1446,7 @@ void cv::ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize
int cn = src.oclchannels();
normalizeAnchor(anchor, ksize);
const static String kernelName = "edgeEnhancingFilter";
const static String kernelName = "adaptiveBilateralFilter";
dst.create(src.size(), src.type());
@ -1480,9 +1495,10 @@ void cv::ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize
//LDATATYPESIZE is sizeof local data store. This is to exemplify effect of LDS on kernel performance
sprintf(build_options,
"-D VAR_PER_CHANNEL=1 -D CALCVAR=1 -D FIXED_WEIGHT=0 -D EXTRA=%d"
"-D VAR_PER_CHANNEL=1 -D CALCVAR=1 -D FIXED_WEIGHT=0 -D EXTRA=%d -D MAX_VAR_VAL=%f -D ABF_GAUSSIAN=%d"
" -D THREADS=%d -D anX=%d -D anY=%d -D ksX=%d -D ksY=%d -D %s",
static_cast<int>(EXTRA), static_cast<int>(blockSizeX), anchor.x, anchor.y, ksize.width, ksize.height, btype);
static_cast<int>(EXTRA), static_cast<float>(maxSigmaColor*maxSigmaColor), static_cast<int>(ABF_GAUSSIAN_ocl),
static_cast<int>(blockSizeX), anchor.x, anchor.y, ksize.width, ksize.height, btype);
std::vector<std::pair<size_t , const void *> > args;
args.push_back(std::make_pair(sizeof(cl_mem), &src.data));

@ -196,9 +196,14 @@ namespace cv
return;
}
if (map1.empty())
map1.swap(map2);
CV_Assert(interpolation == INTER_LINEAR || interpolation == INTER_NEAREST
|| interpolation == INTER_CUBIC || interpolation == INTER_LANCZOS4);
CV_Assert((map1.type() == CV_16SC2 && !map2.data) || (map1.type() == CV_32FC2 && !map2.data) ||
/*|| interpolation == INTER_CUBIC || interpolation == INTER_LANCZOS4*/);
CV_Assert((map1.type() == CV_16SC2 && (map2.empty() || (interpolation == INTER_NEAREST &&
(map2.type() == CV_16UC1 || map2.type() == CV_16SC1)) )) ||
(map1.type() == CV_32FC2 && !map2.data) ||
(map1.type() == CV_32FC1 && map2.type() == CV_32FC1));
CV_Assert(!map2.data || map2.size() == map1.size());
CV_Assert(borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE || borderType == BORDER_WRAP
@ -213,10 +218,14 @@ namespace cv
"BORDER_REFLECT_101", "BORDER_TRANSPARENT" };
String kernelName = "remap";
if ( map1.type() == CV_32FC2 && !map2.data )
if (map1.type() == CV_32FC2 && map2.empty())
kernelName = kernelName + "_32FC2";
else if (map1.type() == CV_16SC2 && !map2.data)
else if (map1.type() == CV_16SC2)
{
kernelName = kernelName + "_16SC2";
if (!map2.empty())
kernelName = kernelName + "_16UC1";
}
else if (map1.type() == CV_32FC1 && map2.type() == CV_32FC1)
kernelName = kernelName + "_2_32FC1";
else
@ -233,16 +242,13 @@ namespace cv
if (interpolation != INTER_NEAREST)
{
int wdepth = std::max(CV_32F, dst.depth());
if (!supportsDouble)
wdepth = std::min(CV_32F, wdepth);
buildOptions = buildOptions
+ format(" -D WT=%s%s -D convertToT=convert_%s%s%s -D convertToWT=convert_%s%s"
" -D convertToWT2=convert_%s2 -D WT2=%s2",
typeMap[wdepth], channelMap[ocn],
typeMap[src.depth()], channelMap[ocn], src.depth() < CV_32F ? "_sat_rte" : "",
typeMap[wdepth], channelMap[ocn],
typeMap[wdepth], typeMap[wdepth]);
+ format(" -D WT=%s%s -D convertToT=convert_%s%s%s -D convertToWT=convert_%s%s"
" -D convertToWT2=convert_%s2 -D WT2=%s2",
typeMap[wdepth], channelMap[ocn],
typeMap[src.depth()], channelMap[ocn], src.depth() < CV_32F ? "_sat_rte" : "",
typeMap[wdepth], channelMap[ocn],
typeMap[wdepth], typeMap[wdepth]);
}
int src_step = src.step / src.elemSize(), src_offset = src.offset / src.elemSize();

@ -160,63 +160,66 @@ static void generateCentersPP(const Mat& _data, Mat& _out_centers,
}
}
void cv::ocl::distanceToCenters(oclMat &dists, oclMat &labels, const oclMat &src, const oclMat &centers, int distType, const oclMat &indices)
void cv::ocl::distanceToCenters(const oclMat &src, const oclMat &centers, Mat &dists, Mat &labels, int distType)
{
CV_Assert(src.cols*src.oclchannels() == centers.cols*centers.oclchannels());
CV_Assert(src.cols * src.channels() == centers.cols * centers.channels());
CV_Assert(src.depth() == CV_32F && centers.depth() == CV_32F);
bool is_label_row_major = false;
ensureSizeIsEnough(1, src.rows, CV_32FC1, dists);
if(labels.empty() || (!labels.empty() && labels.rows == src.rows && labels.cols == 1))
{
ensureSizeIsEnough(src.rows, 1, CV_32SC1, labels);
is_label_row_major = true;
}
CV_Assert(distType == NORM_L1 || distType == NORM_L2SQR);
std::stringstream build_opt_ss;
build_opt_ss
<< (distType == NORM_L1 ? "-D L1_DIST" : "-D L2SQR_DIST")
<< (indices.empty() ? "" : " -D USE_INDEX");
String build_opt = build_opt_ss.str();
dists.create(src.rows, 1, CV_32FC1);
labels.create(src.rows, 1, CV_32SC1);
const int src_step = (int)(src.oclchannels() * src.step / src.elemSize());
const int centers_step = (int)(centers.oclchannels() * centers.step / centers.elemSize());
const int colsNumb = centers.cols*centers.oclchannels();
const int label_step = is_label_row_major ? (int)(labels.step / labels.elemSize()) : 1;
String kernelname = "distanceToCenters";
const int number_of_input = indices.empty() ? src.rows : indices.size().area();
std::stringstream build_opt_ss;
build_opt_ss << (distType == NORM_L1 ? "-D L1_DIST" : "-D L2SQR_DIST");
const int src_offset = (int)src.offset/src.elemSize();
const int centers_offset = (int)centers.offset/centers.elemSize();
int src_step = src.step / src.elemSize1();
int centers_step = centers.step / centers.elemSize1();
int feature_width = centers.cols * centers.oclchannels();
int src_offset = src.offset / src.elemSize1();
int centers_offset = centers.offset / centers.elemSize1();
size_t globalThreads[3] = {number_of_input, 1, 1};
int all_dist_count = src.rows * centers.rows;
oclMat all_dist(1, all_dist_count, CV_32FC1);
std::vector<std::pair<size_t, const void *> > args;
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&src.data));
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&centers.data));
if(!indices.empty())
{
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&indices.data));
}
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&labels.data));
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&dists.data));
args.push_back(std::make_pair(sizeof(cl_int), (void *)&colsNumb));
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&all_dist.data));
args.push_back(std::make_pair(sizeof(cl_int), (void *)&feature_width));
args.push_back(std::make_pair(sizeof(cl_int), (void *)&src_step));
args.push_back(std::make_pair(sizeof(cl_int), (void *)&centers_step));
args.push_back(std::make_pair(sizeof(cl_int), (void *)&label_step));
args.push_back(std::make_pair(sizeof(cl_int), (void *)&number_of_input));
args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.rows));
args.push_back(std::make_pair(sizeof(cl_int), (void *)&centers.rows));
args.push_back(std::make_pair(sizeof(cl_int), (void *)&src_offset));
args.push_back(std::make_pair(sizeof(cl_int), (void *)&centers_offset));
size_t globalThreads[3] = { all_dist_count, 1, 1 };
openCLExecuteKernel(Context::getContext(), &kmeans_kernel,
kernelname, globalThreads, NULL, args, -1, -1, build_opt.c_str());
"distanceToCenters", globalThreads, NULL, args, -1, -1, build_opt_ss.str().c_str());
Mat all_dist_cpu;
all_dist.download(all_dist_cpu);
for (int i = 0; i < src.rows; ++i)
{
Point p;
double minVal;
Rect roi(i * centers.rows, 0, centers.rows, 1);
Mat hdr(all_dist_cpu, roi);
cv::minMaxLoc(hdr, &minVal, NULL, &p);
dists.at<float>(i, 0) = static_cast<float>(minVal);
labels.at<int>(i, 0) = p.x;
}
}
///////////////////////////////////k - means /////////////////////////////////////////////////////////
double cv::ocl::kmeans(const oclMat &_src, int K, oclMat &_bestLabels,
TermCriteria criteria, int attempts, int flags, oclMat &_centers)
{
@ -429,28 +432,19 @@ double cv::ocl::kmeans(const oclMat &_src, int K, oclMat &_bestLabels,
break;
// assign labels
oclMat _dists(1, N, CV_64F);
_bestLabels.upload(_labels);
Mat dists(1, N, CV_64F);
_centers.upload(centers);
distanceToCenters(_src, _centers, dists, _labels);
_bestLabels.upload(_labels);
distanceToCenters(_dists, _bestLabels, _src, _centers);
Mat dists;
_dists.download(dists);
_bestLabels.download(_labels);
float* dist = dists.ptr<float>(0);
compactness = 0;
for( i = 0; i < N; i++ )
{
compactness += (double)dist[i];
}
compactness += (double)dist[i];
}
if( compactness < best_compactness )
{
best_compactness = compactness;
}
}
return best_compactness;

@ -48,22 +48,22 @@
#define T_MEAN_VAR float
#define CONVERT_TYPE convert_uchar_sat
#define F_ZERO (0.0f)
float cvt(uchar val)
inline float cvt(uchar val)
{
return val;
}
float sqr(float val)
inline float sqr(float val)
{
return val * val;
}
float sum(float val)
inline float sum(float val)
{
return val;
}
float clamp1(float var, float learningRate, float diff, float minVar)
static float clamp1(float var, float learningRate, float diff, float minVar)
{
return fmax(var + learningRate * (diff * diff - var), minVar);
}
@ -72,7 +72,7 @@ float clamp1(float var, float learningRate, float diff, float minVar)
#define T_MEAN_VAR float4
#define CONVERT_TYPE convert_uchar4_sat
#define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f)
float4 cvt(const uchar4 val)
inline float4 cvt(const uchar4 val)
{
float4 result;
result.x = val.x;
@ -83,17 +83,17 @@ float4 cvt(const uchar4 val)
return result;
}
float sqr(const float4 val)
inline float sqr(const float4 val)
{
return val.x * val.x + val.y * val.y + val.z * val.z;
}
float sum(const float4 val)
inline float sum(const float4 val)
{
return (val.x + val.y + val.z);
}
float4 clamp1(const float4 var, float learningRate, const float4 diff, float minVar)
static float4 clamp1(const float4 var, float learningRate, const float4 diff, float minVar)
{
float4 result;
result.x = fmax(var.x + learningRate * (diff.x * diff.x - var.x), minVar);
@ -116,14 +116,14 @@ typedef struct
uchar c_shadowVal;
}con_srtuct_t;
void swap(__global float* ptr, int x, int y, int k, int rows, int ptr_step)
static void swap(__global float* ptr, int x, int y, int k, int rows, int ptr_step)
{
float val = ptr[(k * rows + y) * ptr_step + x];
ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x];
ptr[((k + 1) * rows + y) * ptr_step + x] = val;
}
void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_step)
static void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_step)
{
float4 val = ptr[(k * rows + y) * ptr_step + x];
ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x];
@ -412,7 +412,7 @@ __kernel void mog2_kernel(__global T_FRAME * frame, __global int* fgmask, __glob
if (_weight < -prune)
{
_weight = 0.0;
_weight = 0.0f;
nmodes--;
}

@ -85,7 +85,7 @@
#endif
__kernel void
edgeEnhancingFilter_C4_D0(
adaptiveBilateralFilter_C4_D0(
__global const uchar4 * restrict src,
__global uchar4 *dst,
float alpha,
@ -173,14 +173,14 @@ edgeEnhancingFilter_C4_D0(
//find variance of all data
int startLMj;
int endLMj ;
#if CALCVAR
// Top row: don't sum the very last element
for(int extraCnt = 0; extraCnt <=EXTRA; extraCnt++)
{
#if CALCVAR
startLMj = extraCnt;
endLMj = ksY+extraCnt-1;
sumVal =0;
sumValSqr=0;
sumVal = (int4)0;
sumValSqr= (int4)0;
for(int j = startLMj; j < endLMj; j++)
for(int i=-anX; i<=anX; i++)
{
@ -190,9 +190,10 @@ edgeEnhancingFilter_C4_D0(
sumValSqr += mul24(currVal, currVal);
}
var[extraCnt] = convert_float4( ( (sumValSqr * howManyAll)- mul24(sumVal , sumVal) ) ) / ( (float)(howManyAll*howManyAll) ) ;
var[extraCnt] = clamp( convert_float4( ( (sumValSqr * howManyAll)- mul24(sumVal , sumVal) ) ) / ( (float)(howManyAll*howManyAll) ), (float4)(0.1f, 0.1f, 0.1f, 0.1f), (float4)(MAX_VAR_VAL, MAX_VAR_VAL, MAX_VAR_VAL, MAX_VAR_VAL)) ;
#else
var[extraCnt] = (float4)(900.0, 900.0, 900.0, 0.0);
var[extraCnt] = (float4)(MAX_VAR_VAL, MAX_VAR_VAL, MAX_VAR_VAL, MAX_VAR_VAL);
#endif
}
@ -221,32 +222,48 @@ edgeEnhancingFilter_C4_D0(
#else
weight = 1.0f;
#endif
#else
#else // !FIXED_WEIGHT
currVal = convert_int4(data[j][col+anX+i]);
currWRTCenter = currVal-currValCenter;
#if ABF_GAUSSIAN
#if VAR_PER_CHANNEL
weight = var[extraCnt] / (var[extraCnt] + convert_float4(currWRTCenter * currWRTCenter)) *
weight = exp( (float4)(-0.5f, -0.5f, -0.5f, -0.5f) * convert_float4(currWRTCenter * currWRTCenter) / var[extraCnt] )*
(float4)(lut[lut_j*lut_step+anX+i]);
#else
weight = 1.0f/(1.0f+( mul24(currWRTCenter.x, currWRTCenter.x) + mul24(currWRTCenter.y, currWRTCenter.y) +
mul24(currWRTCenter.z, currWRTCenter.z))/(var.x+var.y+var.z));
weight = exp( -0.5f * (mul24(currWRTCenter.x, currWRTCenter.x) + mul24(currWRTCenter.y, currWRTCenter.y) +
mul24(currWRTCenter.z, currWRTCenter.z) ) / (var[extraCnt].x+var[extraCnt].y+var[extraCnt].z) ) * lut[lut_j*lut_step+anX+i];
#endif
#else // !ABF_GAUSSIAN
#if VAR_PER_CHANNEL
weight = var[extraCnt] / (var[extraCnt] + convert_float4(currWRTCenter * currWRTCenter)) *
(float4)(lut[lut_j*lut_step+anX+i]);
#else
weight = ((float)lut[lut_j*lut_step+anX+i]) /(1.0f+( mul24(currWRTCenter.x, currWRTCenter.x) + mul24(currWRTCenter.y, currWRTCenter.y) +
mul24(currWRTCenter.z, currWRTCenter.z))/(var[extraCnt].x+var[extraCnt].y+var[extraCnt].z));
#endif
#endif //ABF_GAUSSIAN
#endif // FIXED_WEIGHT
tmp_sum[extraCnt] += convert_float4(data[j][col+anX+i]) * weight;
totalWeight += weight;
}
}
tmp_sum[extraCnt] /= totalWeight;
if(posX >= 0 && posX < dst_cols && (posY+extraCnt) >= 0 && (posY+extraCnt) < dst_rows)
dst[(dst_startY+extraCnt) * (dst_step>>2)+ dst_startX + col] = convert_uchar4(tmp_sum[extraCnt]);
dst[(dst_startY+extraCnt) * (dst_step>>2)+ dst_startX + col] = convert_uchar4_rtz( (tmp_sum[extraCnt] / (float4)totalWeight) + (float4)0.5f);
#if VAR_PER_CHANNEL
totalWeight = (float4)(0,0,0,0);
#else
totalWeight = 0;
totalWeight = 0.0f;
#endif
}
}
@ -254,7 +271,7 @@ edgeEnhancingFilter_C4_D0(
__kernel void
edgeEnhancingFilter_C1_D0(
adaptiveBilateralFilter_C1_D0(
__global const uchar * restrict src,
__global uchar *dst,
float alpha,
@ -343,10 +360,11 @@ edgeEnhancingFilter_C1_D0(
//find variance of all data
int startLMj;
int endLMj;
#if CALCVAR
// Top row: don't sum the very last element
for(int extraCnt=0; extraCnt<=EXTRA; extraCnt++)
{
#if CALCVAR
startLMj = extraCnt;
endLMj = ksY+extraCnt-1;
sumVal = 0;
@ -361,9 +379,9 @@ edgeEnhancingFilter_C1_D0(
sumValSqr += mul24(currVal, currVal);
}
}
var[extraCnt] = (float)( ( (sumValSqr * howManyAll)- mul24(sumVal , sumVal) ) ) / ( (float)(howManyAll*howManyAll) ) ;
var[extraCnt] = clamp((float)( ( (sumValSqr * howManyAll)- mul24(sumVal , sumVal) ) ) / ( (float)(howManyAll*howManyAll) ) , 0.1f, (float)(MAX_VAR_VAL) );
#else
var[extraCnt] = (float)(900.0);
var[extraCnt] = (float)(MAX_VAR_VAL);
#endif
}
@ -389,19 +407,20 @@ edgeEnhancingFilter_C1_D0(
currVal = (int)(data[j][col+anX+i]) ;
currWRTCenter = currVal-currValCenter;
#if ABF_GAUSSIAN
weight = exp( -0.5f * (float)mul24(currWRTCenter,currWRTCenter)/var[extraCnt]) * lut[lut_j*lut_step+anX+i] ;
#else
weight = var[extraCnt] / (var[extraCnt] + (float)mul24(currWRTCenter,currWRTCenter)) * lut[lut_j*lut_step+anX+i] ;
#endif
#endif
tmp_sum[extraCnt] += (float)(data[j][col+anX+i] * weight);
totalWeight += weight;
}
}
tmp_sum[extraCnt] /= totalWeight;
if(posX >= 0 && posX < dst_cols && (posY+extraCnt) >= 0 && (posY+extraCnt) < dst_rows)
{
dst[(dst_startY+extraCnt) * (dst_step)+ dst_startX + col] = (uchar)(tmp_sum[extraCnt]);
dst[(dst_startY+extraCnt) * (dst_step)+ dst_startX + col] = convert_uchar_rtz(tmp_sum[extraCnt]/totalWeight+0.5f);
}
totalWeight = 0;

@ -292,7 +292,6 @@ __kernel void __attribute__((reqd_work_group_size(8,8,1)))gpuRunHaarClassifierCa
for(int scalei = 0; scalei <loopcount; scalei++)
{
int4 scaleinfo1= info[scalei];
int width = (scaleinfo1.x & 0xffff0000) >> 16;
int height = scaleinfo1.x & 0xffff;
int grpnumperline =(scaleinfo1.y & 0xffff0000) >> 16;
int totalgrp = scaleinfo1.y & 0xffff;

@ -136,8 +136,6 @@ __kernel void gpuRunHaarClassifierCascade_scaled2(
{
int4 scaleinfo1;
scaleinfo1 = info[scalei];
int width = (scaleinfo1.x & 0xffff0000) >> 16;
int height = scaleinfo1.x & 0xffff;
int grpnumperline = (scaleinfo1.y & 0xffff0000) >> 16;
int totalgrp = scaleinfo1.y & 0xffff;
float factor = as_float(scaleinfo1.w);

@ -125,10 +125,12 @@ __kernel void calcHarris(__global const float *Dx, __global const float *Dy, __g
int indexDx = (dx_startY+i)*(dx_step>>2)+(dx_startX+col);
float dx_s = dx_con ? Dx[indexDx] : 0.0f;
dx_data[i] = dx_s;
bool dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows;
int indexDy = (dy_startY+i)*(dy_step>>2)+(dy_startX+col);
float dy_s = dx_con ? Dy[indexDy] : 0.0f;
float dy_s = dy_con ? Dy[indexDy] : 0.0f;
dy_data[i] = dy_s;
data[0][i] = dx_data[i] * dx_data[i];
data[1][i] = dx_data[i] * dy_data[i];
data[2][i] = dy_data[i] * dy_data[i];

@ -124,10 +124,12 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
int indexDx = (dx_startY+i)*(dx_step>>2)+(dx_startX+col);
float dx_s = dx_con ? Dx[indexDx] : 0.0f;
dx_data[i] = dx_s;
bool dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows;
int indexDy = (dy_startY+i)*(dy_step>>2)+(dy_startX+col);
float dy_s = dx_con ? Dy[indexDy] : 0.0f;
float dy_s = dy_con ? Dy[indexDy] : 0.0f;
dy_data[i] = dy_s;
data[0][i] = dx_data[i] * dx_data[i];
data[1][i] = dx_data[i] * dy_data[i];
data[2][i] = dy_data[i] * dy_data[i];

@ -51,6 +51,13 @@
#endif
#endif
enum
{
INTER_BITS = 5,
INTER_TAB_SIZE = 1 << INTER_BITS,
INTER_TAB_SIZE2 = INTER_TAB_SIZE * INTER_TAB_SIZE
};
#ifdef INTER_NEAREST
#define convertToWT
#endif
@ -204,6 +211,36 @@ __kernel void remap_16SC2(__global const T * restrict src, __global T * dst, __g
}
}
__kernel void remap_16SC2_16UC1(__global const T * restrict src, __global T * dst, __global short2 * map1, __global ushort * map2,
int src_offset, int dst_offset, int map1_offset, int map2_offset,
int src_step, int dst_step, int map1_step, int map2_step,
int src_cols, int src_rows, int dst_cols, int dst_rows, T scalar)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < dst_cols && y < dst_rows)
{
int dstIdx = mad24(y, dst_step, x + dst_offset);
int map1Idx = mad24(y, map1_step, x + map1_offset);
int map2Idx = mad24(y, map2_step, x + map2_offset);
int map2Value = convert_int(map2[map2Idx]) & (INTER_TAB_SIZE2 - 1);
int dx = (map2Value & (INTER_TAB_SIZE - 1)) < (INTER_TAB_SIZE >> 1) ? 1 : 0;
int dy = (map2Value >> INTER_BITS) < (INTER_TAB_SIZE >> 1) ? 1 : 0;
int2 gxy = convert_int2(map1[map1Idx]) + (int2)(dx, dy);
int gx = gxy.x, gy = gxy.y;
if (NEED_EXTRAPOLATION(gx, gy))
EXTRAPOLATE(gxy, dst[dstIdx])
else
{
int srcIdx = mad24(gy, src_step, gx + src_offset);
dst[dstIdx] = src[srcIdx];
}
}
}
#elif INTER_LINEAR
__kernel void remap_2_32FC1(__global T const * restrict src, __global T * dst,
@ -229,7 +266,7 @@ __kernel void remap_2_32FC1(__global T const * restrict src, __global T * dst,
int2 map_dataD = (int2)(map_dataA.x + 1, map_dataA.y +1);
float2 _u = map_data - convert_float2(map_dataA);
WT2 u = convertToWT2(convert_int2_rte(convertToWT2(_u) * (WT2)32)) / (WT2)32;
WT2 u = convertToWT2(convert_int2_rte(convertToWT2(_u) * (WT2)INTER_TAB_SIZE)) / (WT2)INTER_TAB_SIZE;
WT scalar = convertToWT(nVal);
WT a = scalar, b = scalar, c = scalar, d = scalar;
@ -282,7 +319,7 @@ __kernel void remap_32FC2(__global T const * restrict src, __global T * dst,
int2 map_dataD = (int2)(map_dataA.x + 1, map_dataA.y + 1);
float2 _u = map_data - convert_float2(map_dataA);
WT2 u = convertToWT2(convert_int2_rte(convertToWT2(_u) * (WT2)32)) / (WT2)32;
WT2 u = convertToWT2(convert_int2_rte(convertToWT2(_u) * (WT2)INTER_TAB_SIZE)) / (WT2)INTER_TAB_SIZE;
WT scalar = convertToWT(nVal);
WT a = scalar, b = scalar, c = scalar, d = scalar;

@ -93,8 +93,8 @@ __kernel void threshold(__global const T * restrict src, int src_offset, int src
#endif
else
{
T array[VECSIZE];
VSTOREN(vecValue, 0, array);
__attribute__(( aligned(sizeof(VT)) )) T array[VECSIZE];
*((VT*)array) = vecValue;
#pragma unroll
for (int i = 0; i < VECSIZE; ++i)
if (gx + i < max_index)

@ -44,81 +44,64 @@
//
//M*/
static float distance_(__global const float * center, __global const float * src, int feature_length)
{
float res = 0;
float4 v0, v1, v2;
int i = 0;
#ifdef L1_DIST
# define DISTANCE(A, B) fabs((A) - (B))
#elif defined L2SQR_DIST
# define DISTANCE(A, B) ((A) - (B)) * ((A) - (B))
#else
# define DISTANCE(A, B) ((A) - (B)) * ((A) - (B))
float4 sum = (float4)(0.0f);
#endif
inline float dist(__global const float * center, __global const float * src, int feature_cols)
{
float res = 0;
float4 tmp4;
int i;
for(i = 0; i < feature_cols / 4; i += 4, center += 4, src += 4)
for ( ; i <= feature_length - 4; i += 4)
{
tmp4 = vload4(0, center) - vload4(0, src);
v0 = vload4(0, center + i);
v1 = vload4(0, src + i);
v2 = v1 - v0;
#ifdef L1_DIST
tmp4 = fabs(tmp4);
v0 = fabs(v2);
sum += v0;
#else
tmp4 *= tmp4;
res += dot(v2, v2);
#endif
res += tmp4.x + tmp4.y + tmp4.z + tmp4.w;
}
for(; i < feature_cols; ++i, ++center, ++src)
#ifdef L1_DIST
res = sum.x + sum.y + sum.z + sum.w;
#endif
for ( ; i < feature_length; ++i)
{
res += DISTANCE(*src, *center);
float t0 = src[i];
float t1 = center[i];
#ifdef L1_DIST
res += fabs(t0 - t1);
#else
float t2 = t0 - t1;
res += t2 * t2;
#endif
}
return res;
}
// to be distinguished with distanceToCenters in kmeans_kernel.cl
__kernel void distanceToCenters(
__global const float *src,
__global const float *centers,
#ifdef USE_INDEX
__global const int *indices,
#endif
__global int *labels,
__global float *dists,
int feature_cols,
int src_step,
int centers_step,
int label_step,
int input_size,
int K,
int offset_src,
int offset_centers
)
__kernel void distanceToCenters(__global const float * src, __global const float * centers,
__global float * dists, int feature_length,
int src_step, int centers_step,
int features_count, int centers_count,
int src_offset, int centers_offset)
{
int gid = get_global_id(0);
float euDist, minval;
int minCentroid;
if(gid >= input_size)
{
return;
}
src += offset_src;
centers += offset_centers;
#ifdef USE_INDEX
src += indices[gid] * src_step;
#else
src += gid * src_step;
#endif
minval = dist(centers, src, feature_cols);
minCentroid = 0;
for(int i = 1 ; i < K; i++)
if (gid < (features_count * centers_count))
{
euDist = dist(centers + i * centers_step, src, feature_cols);
if(euDist < minval)
{
minval = euDist;
minCentroid = i;
}
int feature_index = gid / centers_count;
int center_index = gid % centers_count;
int center_idx = mad24(center_index, centers_step, centers_offset);
int src_idx = mad24(feature_index, src_step, src_offset);
dists[gid] = distance_(centers + center_idx, src + src_idx, feature_length);
}
labels[gid * label_step] = minCentroid;
dists[gid] = minval;
}

@ -69,23 +69,16 @@ __global float* dx, __global float* dy, int dx_step)
}
float bicubicCoeff(float x_)
static float bicubicCoeff(float x_)
{
float x = fabs(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;
}
}
__kernel void warpBackwardKernel(__global const float* I0, int I0_step, int I0_col, int I0_row,
@ -170,12 +163,10 @@ __kernel void warpBackwardKernel(__global const float* I0, int I0_step, int I0_c
}
float readImage(__global const float *image, const int x, const int y, const int rows, const int cols, const int elemCntPerRow)
static float readImage(__global const float *image, const int x, const int y, const int rows, const int cols, const int elemCntPerRow)
{
int i0 = clamp(x, 0, cols - 1);
int j0 = clamp(y, 0, rows - 1);
int i1 = clamp(x + 1, 0, cols - 1);
int j1 = clamp(y + 1, 0, rows - 1);
return image[j0 * elemCntPerRow + i0];
}
@ -303,7 +294,7 @@ __kernel void estimateDualVariablesKernel(__global const float* u1, int u1_col,
}
float divergence(__global const float* v1, __global const float* v2, int y, int x, int v1_step, int v2_step)
static float divergence(__global const float* v1, __global const float* v2, int y, int x, int v1_step, int v2_step)
{
if (x > 0 && y > 0)
@ -407,5 +398,4 @@ __kernel void estimateUKernel(__global const float* I1wx, int I1wx_col, int I1wx
error[y * I1wx_step + x] = n1 + n2;
}
}
}

@ -48,13 +48,8 @@
#include "opencv2/ocl/cl_runtime/cl_runtime.hpp"
#if defined(__GNUC__)
#define openCLSafeCall(expr) ___openCLSafeCall(expr, __FILE__, __LINE__, __func__)
#define openCLVerifyCall(res) ___openCLSafeCall(res, __FILE__, __LINE__, __func__)
#else /* defined(__OPENCLCC__) || defined(__MSVC__) */
#define openCLSafeCall(expr) ___openCLSafeCall(expr, __FILE__, __LINE__)
#define openCLVerifyCall(res) ___openCLSafeCall(res, __FILE__, __LINE__)
#endif
#define openCLSafeCall(expr) ___openCLSafeCall(expr, __FILE__, __LINE__, CV_Func)
#define openCLVerifyCall(res) ___openCLSafeCall(res, __FILE__, __LINE__, CV_Func)
namespace cv

@ -686,9 +686,6 @@ float CvSVM_OCL::predict(const CvMat* samples, CV_OUT CvMat* results) const
}
#else
// TODO fix it
CV_Error(Error::StsNotImplemented, "This part of code contains mistakes. Install AMD BLAS in order to get a correct result or use CPU version of SVM");
double degree1 = 0.0;
if (params.kernel_type == CvSVM::POLY)
degree1 = params.degree;
@ -813,9 +810,6 @@ bool CvSVMSolver_ocl::solve_generic( CvSVMSolutionInfo& si )
}
#else
// TODO fix it
CV_Error(Error::StsNotImplemented, "This part of code contains mistakes. Install AMD BLAS in order to get a correct result or use CPU version of SVM");
double degree1 = 0.0;
if(params->kernel_type == CvSVM::POLY)
degree1 = params->degree;
@ -1000,13 +994,15 @@ void CvSVMKernel_ocl::calc( int vcount, const int row_idx, Qfloat* results, Mat&
//int j;
(this->*calc_func_ocl)( vcount, row_idx, results, src);
// FIXIT #if defined HAVE_CLAMDBLAS
#if !defined(HAVE_CLAMDBLAS)
// nothing
#else
const Qfloat max_val = (Qfloat)(FLT_MAX * 1e-3);
int j;
for( j = 0; j < vcount; j++ )
if( results[j] > max_val )
results[j] = max_val;
// FIXIT #endif
#endif
}
bool CvSVMKernel_ocl::create( const CvSVMParams* _params, Calc_ocl _calc_func, Calc _calc_func1 )
@ -1078,12 +1074,13 @@ void CvSVMKernel_ocl::calc_poly( int vcount, const int row_idx, Qfloat* results,
{
calc_non_rbf_base( vcount, row_idx, results, src);
//FIXIT #if defined HAVE_CLAMDBLAS
#if !defined(HAVE_CLAMDBLAS)
// nothing
#else
CvMat R = cvMat( 1, vcount, QFLOAT_TYPE, results );
if( vcount > 0 )
cvPow( &R, &R, params->degree );
//FIXIT #endif
#endif
}
@ -1091,7 +1088,9 @@ void CvSVMKernel_ocl::calc_sigmoid( int vcount, const int row_idx, Qfloat* resul
{
calc_non_rbf_base( vcount, row_idx, results, src);
// TODO: speedup this
//FIXIT #if defined HAVE_CLAMDBLAS
#if !defined(HAVE_CLAMDBLAS)
// nothing
#else
for(int j = 0; j < vcount; j++ )
{
Qfloat t = results[j];
@ -1101,7 +1100,7 @@ void CvSVMKernel_ocl::calc_sigmoid( int vcount, const int row_idx, Qfloat* resul
else
results[j] = (Qfloat)((e - 1.) / (e + 1.));
}
//FIXIT #endif
#endif
}
CvSVM_OCL::CvSVM_OCL()

@ -42,15 +42,17 @@
#include "test_precomp.hpp"
#define DUMP_INFO_STDOUT(propertyDisplayName, propertyValue) \
#define DUMP_PROPERTY_XML(propertyName, propertyValue) \
do { \
std::cout << (propertyDisplayName) << ": " << (propertyValue) << std::endl; \
std::stringstream ssName, ssValue;\
ssName << propertyName;\
ssValue << propertyValue; \
::testing::Test::RecordProperty(ssName.str(), ssValue.str()); \
} while (false)
#define DUMP_INFO_XML(propertyXMLName, propertyValue) \
#define DUMP_MESSAGE_STDOUT(msg) \
do { \
std::stringstream ss; ss << propertyValue; \
::testing::Test::RecordProperty((propertyXMLName), ss.str()); \
std::cout << msg << std::endl; \
} while (false)
#include "opencv2/ocl/private/opencl_dumpinfo.hpp"

@ -192,13 +192,13 @@ PARAM_TEST_CASE(ArithmTestBase, MatDepth, Channels, bool)
use_roi = GET_PARAM(2);
}
void random_roi()
virtual void random_roi()
{
const int type = CV_MAKE_TYPE(depth, cn);
Size roiSize = randomSize(1, MAX_VALUE);
Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src1, src1_roi, roiSize, srcBorder, type, 2, 11);
Border src1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src1, src1_roi, roiSize, src1Border, type, 2, 11);
Border src2Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src2, src2_roi, roiSize, src2Border, type, -1540, 1740);
@ -214,7 +214,7 @@ PARAM_TEST_CASE(ArithmTestBase, MatDepth, Channels, bool)
cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
generateOclMat(gsrc1_whole, gsrc1_roi, src1, roiSize, srcBorder);
generateOclMat(gsrc1_whole, gsrc1_roi, src1, roiSize, src1Border);
generateOclMat(gsrc2_whole, gsrc2_roi, src2, roiSize, src2Border);
generateOclMat(gdst1_whole, gdst1_roi, dst1, roiSize, dst1Border);
generateOclMat(gdst2_whole, gdst2_roi, dst2, roiSize, dst2Border);
@ -1522,6 +1522,48 @@ OCL_TEST_P(Norm, NORM_L2)
}
}
//// Repeat
struct RepeatTestCase :
public ArithmTestBase
{
int nx, ny;
virtual void random_roi()
{
const int type = CV_MAKE_TYPE(depth, cn);
nx = randomInt(1, 4);
ny = randomInt(1, 4);
Size srcRoiSize = randomSize(1, MAX_VALUE);
Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src1, src1_roi, srcRoiSize, srcBorder, type, 2, 11);
Size dstRoiSize(srcRoiSize.width * nx, srcRoiSize.height * ny);
Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(dst1, dst1_roi, dstRoiSize, dst1Border, type, 5, 16);
generateOclMat(gsrc1_whole, gsrc1_roi, src1, srcRoiSize, srcBorder);
generateOclMat(gdst1_whole, gdst1_roi, dst1, dstRoiSize, dst1Border);
}
};
typedef RepeatTestCase Repeat;
OCL_TEST_P(Repeat, Mat)
{
for (int i = 0; i < LOOP_TIMES; ++i)
{
random_roi();
cv::repeat(src1_roi, ny, nx, dst1_roi);
cv::ocl::repeat(gsrc1_roi, ny, nx, gdst1_roi);
Near();
}
}
//////////////////////////////////////// Instantiation /////////////////////////////////////////
INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool(), Bool()));
@ -1557,5 +1599,6 @@ INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, Combine(Values(CV_8U, CV_8S, CV_16U
INSTANTIATE_TEST_CASE_P(Arithm, SetIdentity, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
INSTANTIATE_TEST_CASE_P(Arithm, MeanStdDev, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
INSTANTIATE_TEST_CASE_P(Arithm, Norm, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
INSTANTIATE_TEST_CASE_P(Arithm, Repeat, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
#endif // HAVE_OPENCL

@ -338,8 +338,8 @@ OCL_TEST_P(AdaptiveBilateral, Mat)
{
random_roi();
adaptiveBilateralFilter(src_roi, dst_roi, kernelSize, 5, Point(-1, -1), borderType); // TODO anchor
ocl::adaptiveBilateralFilter(gsrc_roi, gdst_roi, kernelSize, 5, Point(-1, -1), borderType);
adaptiveBilateralFilter(src_roi, dst_roi, kernelSize, 5, 1, Point(-1, -1), borderType); // TODO anchor
ocl::adaptiveBilateralFilter(gsrc_roi, gdst_roi, kernelSize, 5, 1, Point(-1, -1), borderType);
Near();
}

@ -142,18 +142,23 @@ PARAM_TEST_CASE(CopyMakeBorder, MatDepth, // depth
void random_roi()
{
border = randomBorder(0, MAX_VALUE << 2);
val = randomScalar(-MAX_VALUE, MAX_VALUE);
Size roiSize = randomSize(1, MAX_VALUE);
Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
randomSubMat(src, src_roi, roiSize, srcBorder, type, 5, 256);
randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
randomSubMat(dst_whole, dst_roi, roiSize, dstBorder, type, 5, 16);
dstBorder.top += border.top;
dstBorder.lef += border.lef;
dstBorder.rig += border.rig;
dstBorder.bot += border.bot;
randomSubMat(dst_whole, dst_roi, roiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE);
generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder);
generateOclMat(gdst_whole, gdst_roi, dst_whole, roiSize, dstBorder);
border = randomBorder(0, MAX_VALUE << 2);
val = randomScalar(-MAX_VALUE, MAX_VALUE);
}
void Near(double threshold = 0.0)
@ -207,11 +212,19 @@ struct CornerTestBase :
Mat image = readImageType("gpu/stereobm/aloe-L.png", type);
ASSERT_FALSE(image.empty());
bool isFP = CV_MAT_DEPTH(type) >= CV_32F;
float val = 255.0f;
if (isFP)
{
image.convertTo(image, -1, 1.0 / 255);
val /= 255.0f;
}
Size roiSize = image.size();
Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
Size wholeSize = Size(roiSize.width + srcBorder.lef + srcBorder.rig, roiSize.height + srcBorder.top + srcBorder.bot);
src = randomMat(wholeSize, type, -255, 255, false);
src = randomMat(wholeSize, type, -val, val, false);
src_roi = src(Rect(srcBorder.lef, srcBorder.top, roiSize.width, roiSize.height));
image.copyTo(src_roi);
@ -522,7 +535,7 @@ INSTANTIATE_TEST_CASE_P(Imgproc, CornerMinEigenVal, Combine(
Bool()));
INSTANTIATE_TEST_CASE_P(Imgproc, CornerHarris, Combine(
Values((MatType)CV_8UC1), // TODO does not work properly with CV_32FC1
Values((MatType)CV_8UC1, CV_32FC1),
Values(3, 5),
Values( (int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, (int)BORDER_REFLECT_101),
Bool()));
@ -567,14 +580,11 @@ INSTANTIATE_TEST_CASE_P(Imgproc, ColumnSum, Combine(
Bool()));
INSTANTIATE_TEST_CASE_P(ImgprocTestBase, CopyMakeBorder, Combine(
testing::Range((MatDepth)CV_8U, (MatDepth)CV_USRTYPE1),
testing::Values((Channels)1, (Channels)4),
testing::Values((MatDepth)CV_8U, (MatDepth)CV_16S, (MatDepth)CV_32S, (MatDepth)CV_32F),
testing::Values(Channels(1), Channels(3), (Channels)4),
Bool(), // border isolated or not
Values((Border)BORDER_CONSTANT,
(Border)BORDER_REPLICATE,
(Border)BORDER_REFLECT,
(Border)BORDER_WRAP,
(Border)BORDER_REFLECT_101),
Values((Border)BORDER_REPLICATE, (Border)BORDER_REFLECT,
(Border)BORDER_WRAP, (Border)BORDER_REFLECT_101),
Bool()));
#endif // HAVE_OPENCL

@ -61,7 +61,7 @@ PARAM_TEST_CASE(Kmeans, int, int, int)
int type;
int K;
int flags;
cv::Mat src ;
Mat src ;
ocl::oclMat d_src, d_dists;
Mat labels, centers;
@ -73,7 +73,7 @@ PARAM_TEST_CASE(Kmeans, int, int, int)
flags = GET_PARAM(2);
// MWIDTH=256, MHEIGHT=256. defined in utility.hpp
cv::Size size = cv::Size(MWIDTH, MHEIGHT);
Size size = Size(MWIDTH, MHEIGHT);
src.create(size, type);
int row_idx = 0;
const int max_neighbour = MHEIGHT / K - 1;
@ -159,15 +159,15 @@ INSTANTIATE_TEST_CASE_P(OCL_ML, Kmeans, Combine(
/////////////////////////////// DistanceToCenters //////////////////////////////////////////
CV_ENUM(DistType, NORM_L1, NORM_L2SQR);
CV_ENUM(DistType, NORM_L1, NORM_L2SQR)
PARAM_TEST_CASE(distanceToCenters, DistType, bool)
{
cv::Size size;
int distType;
bool useRoi;
cv::Mat src, centers, src_roi, centers_roi;
cv::ocl::oclMat ocl_src, ocl_centers, ocl_src_roi, ocl_centers_roi;
Mat src, centers, src_roi, centers_roi;
ocl::oclMat ocl_src, ocl_centers, ocl_src_roi, ocl_centers_roi;
virtual void SetUp()
{
@ -177,70 +177,59 @@ PARAM_TEST_CASE(distanceToCenters, DistType, bool)
void random_roi()
{
Size roiSize_src = randomSize(10,1000);
Size roiSize_centers = randomSize(10, 1000);
roiSize_src.width = roiSize_centers.width;
Size roiSizeSrc = randomSize(1, MAX_VALUE);
Size roiSizeCenters = randomSize(1, MAX_VALUE);
roiSizeSrc.width = roiSizeCenters.width;
Border srcBorder = randomBorder(0, useRoi ? 500 : 0);
randomSubMat(src, src_roi, roiSize_src, srcBorder, CV_32FC1, -SHRT_MAX, SHRT_MAX);
Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
randomSubMat(src, src_roi, roiSizeSrc, srcBorder, CV_32FC1, -MAX_VALUE, MAX_VALUE);
Border centersBorder = randomBorder(0, useRoi ? 500 : 0);
randomSubMat(centers, centers_roi, roiSize_centers, centersBorder, CV_32FC1, -SHRT_MAX, SHRT_MAX);
for(int i = 0; i<centers.rows; i++)
centers.at<float>(i, randomInt(0,centers.cols-1)) = (float)randomDouble(SHRT_MAX, INT_MAX);
randomSubMat(centers, centers_roi, roiSizeCenters, centersBorder, CV_32FC1, -MAX_VALUE, MAX_VALUE);
generateOclMat(ocl_src, ocl_src_roi, src, roiSize_src, srcBorder);
generateOclMat(ocl_centers, ocl_centers_roi, centers, roiSize_centers, centersBorder);
for (int i = 0; i < centers.rows; i++)
centers.at<float>(i, randomInt(0, centers.cols)) = (float)randomDouble(SHRT_MAX, INT_MAX);
generateOclMat(ocl_src, ocl_src_roi, src, roiSizeSrc, srcBorder);
generateOclMat(ocl_centers, ocl_centers_roi, centers, roiSizeCenters, centersBorder);
}
};
OCL_TEST_P(distanceToCenters, Accuracy)
{
for(int j = 0; j< LOOP_TIMES; j++)
for (int j = 0; j < LOOP_TIMES; j++)
{
random_roi();
cv::ocl::oclMat ocl_dists;
cv::ocl::oclMat ocl_labels;
cv::ocl::distanceToCenters(ocl_dists,ocl_labels,ocl_src_roi, ocl_centers_roi, distType);
Mat labels, dists;
ocl_labels.download(labels);
ocl_dists.download(dists);
ocl::distanceToCenters(ocl_src_roi, ocl_centers_roi, dists, labels, distType);
ASSERT_EQ(ocl_dists.cols, ocl_labels.rows);
EXPECT_EQ(dists.size(), labels.size());
Mat batch_dists;
cv::batchDistance(src_roi, centers_roi, batch_dists, CV_32FC1, noArray(), distType);
std::vector<double> gold_dists_v;
std::vector<float> gold_dists_v;
gold_dists_v.reserve(batch_dists.rows);
for(int i = 0; i<batch_dists.rows; i++)
for (int i = 0; i < batch_dists.rows; i++)
{
Mat r = batch_dists.row(i);
double mVal;
Point mLoc;
minMaxLoc(r, &mVal, NULL, &mLoc, NULL);
int ocl_label = *(int*)labels.row(i).col(0).data;
ASSERT_EQ(mLoc.x, ocl_label);
int ocl_label = labels.at<int>(i, 0);
EXPECT_EQ(mLoc.x, ocl_label);
gold_dists_v.push_back(mVal);
gold_dists_v.push_back(static_cast<float>(mVal));
}
Mat gold_dists(gold_dists_v);
dists.convertTo(dists, CV_64FC1);
double relative_error = cv::norm(gold_dists.t(), dists, NORM_INF|NORM_RELATIVE);
double relative_error = cv::norm(Mat(gold_dists_v), dists, NORM_INF | NORM_RELATIVE);
ASSERT_LE(relative_error, 1e-5);
}
}
INSTANTIATE_TEST_CASE_P (OCL_ML, distanceToCenters, Combine(DistType::all(), Bool()) );
INSTANTIATE_TEST_CASE_P (OCL_ML, distanceToCenters, Combine(DistType::all(), Bool()));
#endif

@ -126,8 +126,6 @@ OCL_TEST_P(KNN, Accuracy)
INSTANTIATE_TEST_CASE_P(OCL_ML, KNN, Combine(Values(6, 5), Values(Size(200, 400), Size(300, 600)),
Values(4, 3), Values(false, true)));
#ifdef HAVE_CLAMDBLAS // TODO does not work non-blas version of SVM
////////////////////////////////SVM/////////////////////////////////////////////////
PARAM_TEST_CASE(SVM_OCL, int, int, int)
@ -308,6 +306,4 @@ INSTANTIATE_TEST_CASE_P(OCL_ML, SVM_OCL, testing::Combine(
Values(2, 3, 4)
));
#endif // HAVE_CLAMDBLAS
#endif // HAVE_OPENCL

@ -210,8 +210,9 @@ PARAM_TEST_CASE(Haar, int, CascadeName)
OCL_TEST_P(Haar, FaceDetect)
{
cascade.detectMultiScale(d_img, oclfaces, 1.1, 3,
flags, Size(30, 30));
cascade.detectMultiScale(d_img, oclfaces, 1.1, 3,
flags,
Size(30, 30), Size(0, 0));
cpucascade.detectMultiScale(img, faces, 1.1, 3,
flags,

@ -355,6 +355,7 @@ INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_NEAREST, Combine(
Values(1, 2, 3, 4),
Values(pair<MatType, MatType>((MatType)CV_32FC1, (MatType)CV_32FC1),
pair<MatType, MatType>((MatType)CV_32FC2, noType),
pair<MatType, MatType>((MatType)CV_16SC2, (MatType)CV_16UC1),
pair<MatType, MatType>((MatType)CV_16SC2, noType)),
Values((Border)BORDER_CONSTANT,
(Border)BORDER_REPLICATE,

@ -66,7 +66,9 @@ foreach(m ${OPENCV_MODULE_${the_module}_DEPS})
endif()
endif()
string(REPLACE ".." "__" srcname "${srcname}")
get_filename_component(srcname_we ${srcname} NAME_WE)
#NAME_WE intentionally not used since it interprets first period as start of extension (http://cmake.org/Bug/view.php?id=12282)
get_filename_component(srcname_we "${srcname}" NAME)
string(REGEX REPLACE "\\.[^.]+$" "" srcname_we "${srcname_we}")
string(REGEX REPLACE <SRC_NAME_WE> "${srcname_we}" objpath2 "${objpath1}")
string(REGEX REPLACE <RELATIVE_SRC_NAME> "${srcname}" objpath3 "${objpath2}")
endif()
@ -92,7 +94,9 @@ macro(ios_include_3party_libs)
endif()
string(REPLACE ".." "__" srcname "${srcname}")
get_filename_component(srcname_we ${srcname} NAME_WE)
#NAME_WE intentionally not used since it interprets first period as start of extension (http://cmake.org/Bug/view.php?id=12282)
get_filename_component(srcname_we "${srcname}" NAME)
string(REGEX REPLACE "\\.[^.]+$" "" srcname_we "${srcname_we}")
string(REGEX REPLACE <SRC_NAME_WE> "${srcname_we}" objpath2 "${objpath1}")
string(REGEX REPLACE <RELATIVE_SRC_NAME> "${srcname}" objpath3 "${objpath2}")

@ -1,8 +1,8 @@
<?xml version="1.0" encoding="utf-8"?>
<manifest xmlns:android="http://schemas.android.com/apk/res/android"
package="org.opencv.engine"
android:versionCode="213@ANDROID_PLATFORM_VERSION_CODE@"
android:versionName="2.13" >
android:versionCode="214@ANDROID_PLATFORM_VERSION_CODE@"
android:versionName="2.14" >
<uses-sdk android:minSdkVersion="@ANDROID_NATIVE_API_LEVEL@" />
<uses-feature android:name="android.hardware.touchscreen" android:required="false"/>

@ -14,20 +14,20 @@ manually using adb tool:
.. code-block:: sh
adb install OpenCV-2.4.6-android-sdk/apk/OpenCV_2.4.6_Manager_2.9_<platform>.apk
adb install OpenCV-2.4.7-android-sdk/apk/OpenCV_2.4.7_Manager_2.14_<platform>.apk
Use the table below to determine proper OpenCV Manager package for your device:
+------------------------------+--------------+---------------------------------------------------+
| Hardware Platform | Android ver. | Package name |
+==============================+==============+===================================================+
| armeabi-v7a (ARMv7-A + NEON) | >= 2.3 | OpenCV_2.4.6_Manager_2.9_armv7a-neon.apk |
+------------------------------+--------------+---------------------------------------------------+
| armeabi-v7a (ARMv7-A + NEON) | = 2.2 | OpenCV_2.4.6_Manager_2.9_armv7a-neon-android8.apk |
+------------------------------+--------------+---------------------------------------------------+
| armeabi (ARMv5, ARMv6) | >= 2.3 | OpenCV_2.4.6_Manager_2.9_armeabi.apk |
+------------------------------+--------------+---------------------------------------------------+
| Intel x86 | >= 2.3 | OpenCV_2.4.6_Manager_2.9_x86.apk |
+------------------------------+--------------+---------------------------------------------------+
| MIPS | >= 2.3 | OpenCV_2.4.6_Manager_2.9_mips.apk |
+------------------------------+--------------+---------------------------------------------------+
+------------------------------+--------------+----------------------------------------------------+
| Hardware Platform | Android ver. | Package name |
+==============================+==============+====================================================+
| armeabi-v7a (ARMv7-A + NEON) | >= 2.3 | OpenCV_2.4.7_Manager_2.14_armv7a-neon.apk |
+------------------------------+--------------+----------------------------------------------------+
| armeabi-v7a (ARMv7-A + NEON) | = 2.2 | OpenCV_2.4.7_Manager_2.14_armv7a-neon-android8.apk |
+------------------------------+--------------+----------------------------------------------------+
| armeabi (ARMv5, ARMv6) | >= 2.3 | OpenCV_2.4.7_Manager_2.14_armeabi.apk |
+------------------------------+--------------+----------------------------------------------------+
| Intel x86 | >= 2.3 | OpenCV_2.4.7_Manager_2.14_x86.apk |
+------------------------------+--------------+----------------------------------------------------+
| MIPS | >= 2.3 | OpenCV_2.4.7_Manager_2.14_mips.apk |
+------------------------------+--------------+----------------------------------------------------+

@ -13,7 +13,9 @@ int main( int argc, const char** argv )
{
const char* keys =
"{ i input | | specify input image }"
"{ k ksize | 5 | specify kernel size }"
"{ k ksize | 11 | specify kernel size }"
"{ s sSpace | 3 | specify sigma space }"
"{ c sColor | 30 | specify max color }"
"{ h help | false | print help message }";
CommandLineParser cmd(argc, argv, keys);
@ -27,27 +29,36 @@ int main( int argc, const char** argv )
string src_path = cmd.get<string>("i");
int ks = cmd.get<int>("k");
const char * winName[] = {"input", "adaptive bilateral CPU", "adaptive bilateral OpenCL", "bilateralFilter OpenCL"};
const char * winName[] = {"input", "ABF OpenCL", "BF OpenCL"};
Mat src = imread(src_path), abFilterCPU;
Mat src = imread(src_path);
if (src.empty())
{
cout << "error read image: " << src_path << endl;
return EXIT_FAILURE;
}
ocl::oclMat dsrc(src), dABFilter, dBFilter;
double sigmaSpace = cmd.get<int>("s");
// sigma for checking pixel values. This is used as is in the "normal" bilateral filter,
// and it is used as an upper clamp on the adaptive case.
double sigmacolor = cmd.get<int>("c");
ocl::oclMat dsrc(src), dABFilter, dBFilter;
Size ksize(ks, ks);
adaptiveBilateralFilter(src,abFilterCPU, ksize, 10);
ocl::adaptiveBilateralFilter(dsrc, dABFilter, ksize, 10);
ocl::bilateralFilter(dsrc, dBFilter, ks, 30, 9);
// ksize is the total width/height of neighborhood used to calculate local variance.
// sigmaSpace is not a priori related to ksize/2.
ocl::adaptiveBilateralFilter(dsrc, dABFilter, ksize, sigmaSpace, sigmacolor);
ocl::bilateralFilter(dsrc, dBFilter, ks, sigmacolor, sigmaSpace);
Mat abFilter = dABFilter, bFilter = dBFilter;
ocl::finish();
imshow(winName[0], src);
imshow(winName[1], abFilterCPU);
imshow(winName[2], abFilter);
imshow(winName[3], bFilter);
imshow(winName[1], abFilter);
imshow(winName[2], bFilter);
waitKey();
return EXIT_SUCCESS;

Loading…
Cancel
Save