added tests

modules/ocl/include/opencv2/ocl.hpp

@@ -1558,10 +1558,12 @@ namespace cv
             /*
              * Compute the descriptors for a set of keypoints in an image.
              * image        The image.
-             * keypoints    The input keypoints. Keypoints for which a descriptor cannot be computed are removed.
+             * keypoints    The input keypoints.
              * descriptors  Copmputed descriptors. Row i is the descriptor for keypoint i.
              */
-            void compute( const oclMat& image, oclMat& keypoints, oclMat& descriptors ) const;
+            void compute( const oclMat& image, const oclMat& keypoints, oclMat& mask, oclMat& descriptors ) const;
+
+            static int getBorderSize();
         protected:
 
             int bytes;

modules/ocl/perf/perf_brief.cpp

@@ -50,54 +50,64 @@ using namespace ocl;
 using namespace perf;
 
 ///////////// BRIEF ////////////////////////
+typedef TestBaseWithParam<std::tr1::tuple<std::string, int, size_t> > OCL_BRIEF;
 
-typedef TestBaseWithParam<std::tr1::tuple<std::string, int> > OCL_BRIEF;
+PERF_TEST_P( OCL_BRIEF, extract, testing::Combine(
-
+                                                   testing::Values( string( "gpu/opticalflow/rubberwhale1.png" ),
-#define BRIEF_IMAGES \
+                                                                    string( "gpu/stereobm/aloe-L.png" )
-    "cv/detectors_descriptors_evaluation/images_datasets/leuven/img1.png",\
+                                                                    ), testing::Values( 16, 32, 64 ), testing::Values( 250, 500, 1000, 2500, 3000 ) ) )
-    "stitching/a3.png"
-
-PERF_TEST_P( OCL_BRIEF, extract, testing::Combine( testing::Values( BRIEF_IMAGES ), testing::Values( 16, 32, 64 ) ) )
 {
-    const int threshold = 20;
     const std::string filename = std::tr1::get<0>(GetParam( ));
     const int bytes = std::tr1::get<1>(GetParam( ));
-    const Mat img = imread( getDataPath( filename ), IMREAD_GRAYSCALE );
+    const size_t numKp = std::tr1::get<2>(GetParam( ));
-    ASSERT_FALSE( img.empty( ) );
+
+    Mat img = imread( getDataPath( filename ), IMREAD_GRAYSCALE );
+    ASSERT_TRUE( !img.empty( ) ) << "no input image";
+
+    int threshold = 15;
+    std::vector<KeyPoint> keypoints;
+    while (threshold > 0 && keypoints.size( ) < numKp)
+    {
+        FastFeatureDetector fast( threshold );
+        fast.detect( img, keypoints, Mat( ) );
+        threshold -= 5;
+        KeyPointsFilter::runByImageBorder( keypoints, img.size( ), BRIEF_OCL::getBorderSize( ) );
+    }
+    ASSERT_TRUE( keypoints.size( ) >= numKp ) << "not enough keypoints";
+    keypoints.resize( numKp );
 
     if ( RUN_OCL_IMPL )
     {
-        oclMat d_img( img );
+        Mat kpMat( 2, keypoints.size( ), CV_32FC1 );
-        oclMat d_keypoints;
+        for ( size_t i = 0; i < keypoints.size( ); ++i )
-        FAST_OCL fast( threshold );
-        fast( d_img, oclMat( ), d_keypoints );
-
-        BRIEF_OCL brief( bytes );
-
-        OCL_TEST_CYCLE( )
         {
-            oclMat d_descriptors;
+            kpMat.col( i ).row( 0 ) = keypoints[i].pt.x;
-            brief.compute( d_img, d_keypoints, d_descriptors );
+            kpMat.col( i ).row( 1 ) = keypoints[i].pt.y;
         }
-
+        BRIEF_OCL brief( bytes );
-        std::vector<KeyPoint> ocl_keypoints;
+        oclMat imgCL( img ), keypointsCL(kpMat), mask;
-        fast.downloadKeypoints( d_keypoints, ocl_keypoints );
+        while (next( ))
-        SANITY_CHECK_KEYPOINTS( ocl_keypoints );
+        {
+            startTimer( );
+            oclMat descriptorsCL;
+            brief.compute( imgCL, keypointsCL, mask, descriptorsCL );
+            cv::ocl::finish( );
+            stopTimer( );
+        }
+        SANITY_CHECK_NOTHING( )
     }
     else if ( RUN_PLAIN_IMPL )
     {
-        std::vector<KeyPoint> keypoints;
-        FAST( img, keypoints, threshold );
-
         BriefDescriptorExtractor brief( bytes );
 
-        TEST_CYCLE( )
+        while (next( ))
         {
+            startTimer( );
             Mat descriptors;
             brief.compute( img, keypoints, descriptors );
+            stopTimer( );
         }
-
+        SANITY_CHECK_NOTHING( )
-        SANITY_CHECK_KEYPOINTS( keypoints );
     }
     else
         OCL_PERF_ELSE;

modules/ocl/src/brief.cpp

@@ -53,35 +53,39 @@ BRIEF_OCL::BRIEF_OCL( int _bytes ) : bytes( _bytes )
 {
 }
 
-void
+void BRIEF_OCL::compute( const oclMat& image, const oclMat& keypoints, oclMat& mask, oclMat& descriptors ) const
-BRIEF_OCL::compute( const oclMat& image, oclMat& keypoints, oclMat& descriptors ) const
 {
-    oclMat grayImage = image;
+    CV_Assert( image.type( ) == CV_8UC1 );
-    if ( image.type( ) != CV_8U ) cvtColor( image, grayImage, COLOR_BGR2GRAY );
+    if ( keypoints.size( ).area( ) == 0 ) return;
-
+    descriptors = oclMat( Mat( keypoints.cols, bytes, CV_8UC1 ) );
+    if( mask.cols != keypoints.cols )
+    {
+        mask = oclMat( Mat::ones( 1, keypoints.cols, CV_8UC1 ) );
+    }
     oclMat sum;
-    integral( grayImage, sum, CV_32S );
+    integral( image, sum, CV_32S );
     cl_mem sumTexture = bindTexture( sum );
-
-    //TODO filter keypoints by border
-
-    descriptors = oclMat( keypoints.cols, bytes, CV_8U );
-
     std::stringstream build_opt;
-    build_opt << " -D BYTES=" << bytes << " -D KERNEL_SIZE=" << KERNEL_SIZE;
+    build_opt
-
+            << " -D BYTES=" << bytes
+            << " -D KERNEL_SIZE=" << KERNEL_SIZE
+            << " -D BORDER=" << getBorderSize();
     const String kernelname = "extractBriefDescriptors";
     size_t localThreads[3] = {bytes, 1, 1};
     size_t globalThreads[3] = {keypoints.cols * bytes, 1, 1};
-
+    Context* ctx = Context::getContext( );
     std::vector< std::pair<size_t, const void *> > args;
     args.push_back( std::make_pair( sizeof (cl_mem), (void *) &sumTexture ) );
     args.push_back( std::make_pair( sizeof (cl_mem), (void *) &keypoints.data ) );
     args.push_back( std::make_pair( sizeof (cl_int), (void *) &keypoints.step ) );
     args.push_back( std::make_pair( sizeof (cl_mem), (void *) &descriptors.data ) );
     args.push_back( std::make_pair( sizeof (cl_int), (void *) &descriptors.step ) );
-
+    args.push_back( std::make_pair( sizeof (cl_mem), (void *) &mask.data ) );
-    Context* ctx = Context::getContext( );
     openCLExecuteKernel( ctx, &brief, kernelname, globalThreads, localThreads, args, -1, -1, build_opt.str( ).c_str( ) );
     openCLFree( sumTexture );
 }
+
+int BRIEF_OCL::getBorderSize( )
+{
+    return PATCH_SIZE / 2 + KERNEL_SIZE / 2;
+}

modules/ocl/src/opencl/brief.cl

@@ -41,15 +41,16 @@
 //
 //M*/
 
-#define X_ROW 0
-#define Y_ROW 1
-
 #ifndef BYTES
  #define BYTES 16
 #endif
 
 #ifndef KERNEL_SIZE
- #define KERNEL_SIZE 32
+ #define KERNEL_SIZE 9
+#endif
+
+#ifndef BORDER
+ #define BORDER 0
 #endif
 
 #define HALF_KERNEL (KERNEL_SIZE/2)
@@ -128,29 +129,45 @@ __constant char tests[32 * BYTES] =
 #endif
 };
 
-inline int smoothedSum(__read_only image2d_t sum, const int2 pt)
+inline int smoothedSum(__read_only image2d_t sum, const int2 kpPos, const int2 pt)
 {
-    return ( read_imagei( sum, sampler, pt + (int2)(  HALF_KERNEL + 1,  HALF_KERNEL + 1 ))
+    return ( read_imagei( sum, sampler, kpPos + pt + (int2)(  HALF_KERNEL + 1,  HALF_KERNEL + 1 ))
-           - read_imagei( sum, sampler, pt + (int2)( -HALF_KERNEL,      HALF_KERNEL + 1 ))
+           - read_imagei( sum, sampler, kpPos + pt + (int2)( -HALF_KERNEL,      HALF_KERNEL + 1 ))
-           - read_imagei( sum, sampler, pt + (int2)(  HALF_KERNEL + 1, -HALF_KERNEL ))
+           - read_imagei( sum, sampler, kpPos + pt + (int2)(  HALF_KERNEL + 1, -HALF_KERNEL ))
-           + read_imagei( sum, sampler, pt + (int2)( -HALF_KERNEL,     -HALF_KERNEL ))).x;
+           + read_imagei( sum, sampler, kpPos + pt + (int2)( -HALF_KERNEL,     -HALF_KERNEL ))).x;
 }
 
 __kernel void extractBriefDescriptors(
-    __read_only image2d_t sumImg, __global float* keypoints, int kpRowStep, __global uchar* descriptors, int dscRowStep)
+    __read_only image2d_t sumImg,
+    __global float* keypoints, int kpRowStep,
+    __global uchar* descriptors, int dscRowStep,
+    __global uchar* mask)
 {
     const int  byte  = get_local_id(0);
     const int  kpId  = get_group_id(0);
-    const int2 kpPos = (int2)(keypoints[X_ROW * (kpRowStep/4) + kpId] + 0.5, keypoints[Y_ROW * (kpRowStep/4) + kpId] + 0.5);
 
+    if( !mask[kpId])
+    {
+        return;
+    }
+    const float2 kpPos = (float2)(keypoints[kpId], keypoints[kpRowStep/4 + kpId]);
+    if( kpPos.x < BORDER
+     || kpPos.y < BORDER
+     || kpPos.x >= (get_image_width(  sumImg ) - BORDER)
+     || kpPos.y >= (get_image_height( sumImg ) - BORDER) )
+    {
+        if( byte == 0) mask[kpId] = 0;
+        return;
+    }
     uchar descByte = 0;
+    const int2 pt = (int2)( kpPos.x + 0.5f, kpPos.y + 0.5f );
     for(int i = 0; i<8; ++i)
     {
         descByte |= (
-            smoothedSum(sumImg, (int2)( tests[byte * 32 + (i * 4) + 0], tests[byte * 32 + (i * 4) + 1] ))
+            smoothedSum(sumImg, pt, (int2)( tests[byte * 32 + (i * 4) + 1], tests[byte * 32 + (i * 4) + 0] ))
-          < smoothedSum(sumImg, (int2)( tests[byte * 32 + (i * 4) + 2], tests[byte * 32 + (i * 4) + 3] ))
+          < smoothedSum(sumImg, pt, (int2)( tests[byte * 32 + (i * 4) + 3], tests[byte * 32 + (i * 4) + 2] ))
           ) << (7-i);
     }
-
     descriptors[kpId * dscRowStep + byte] = descByte;
+    if( byte == 0) mask[kpId] = 1;
 }

modules/ocl/test/test_brief.cpp

@@ -0,0 +1,115 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009-2010, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Matthias Bady aegirxx ==> gmail.com
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "test_precomp.hpp"
+
+using namespace std;
+using namespace cv;
+using namespace ocl;
+
+#ifdef HAVE_OPENCL
+
+namespace
+{
+IMPLEMENT_PARAM_CLASS( BRIEF_Bytes, int )
+}
+
+PARAM_TEST_CASE( BRIEF, BRIEF_Bytes )
+{
+    int bytes;
+
+    virtual void SetUp( )
+    {
+        bytes = GET_PARAM( 0 );
+    }
+};
+
+OCL_TEST_P( BRIEF, Accuracy )
+{
+    Mat img = readImage( "gpu/opticalflow/rubberwhale1.png", IMREAD_GRAYSCALE );
+    ASSERT_TRUE( !img.empty( ) ) << "no input image";
+
+    FastFeatureDetector fast( 20 );
+    std::vector<KeyPoint> keypoints;
+    fast.detect( img, keypoints, Mat( ) );
+
+    Mat descriptorsGold;
+    BriefDescriptorExtractor brief( bytes );
+    brief.compute( img, keypoints, descriptorsGold );
+
+    Mat kpMat( 2, keypoints.size( ), CV_32FC1 );
+    for ( size_t i = 0; i < keypoints.size( ); ++i )
+    {
+        kpMat.col( i ).row( 0 ) = keypoints[i].pt.x;
+        kpMat.col( i ).row( 1 ) = keypoints[i].pt.y;
+    }
+    oclMat imgOcl( img ), keypointsOcl( kpMat ), descriptorsOcl, maskOcl;
+
+    BRIEF_OCL briefOcl( bytes );
+    briefOcl.compute( imgOcl, keypointsOcl, maskOcl, descriptorsOcl );
+    Mat mask, descriptors;
+    maskOcl.download( mask );
+    descriptorsOcl.download( descriptors );
+
+    const int numDesc = cv::countNonZero( mask );
+    if ( numDesc != descriptors.cols )
+    {
+        size_t idx = 0;
+        Mat tmp( numDesc, bytes, CV_8UC1 );
+        for ( int i = 0; i < descriptors.rows; ++i )
+        {
+            if ( mask.at<uchar>(i) )
+            {
+                descriptors.row( i ).copyTo( tmp.row( idx++ ) );
+            }
+        }
+        descriptors = tmp;
+    }
+    ASSERT_TRUE( descriptors.size( ) == descriptorsGold.size( ) ) << "Different number of descriptors";
+    ASSERT_TRUE( 0 == norm( descriptors, descriptorsGold, NORM_HAMMING ) ) << "Descriptors different";
+}
+
+INSTANTIATE_TEST_CASE_P( OCL_Features2D, BRIEF, testing::Values( 16, 32, 64 ) );
+#endif