/*
* 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;
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;
-
-#define BRIEF_IMAGES \
- "cv/detectors_descriptors_evaluation/images_datasets/leuven/img1.png",\
- "stitching/a3.png"
-
-PERF_TEST_P( OCL_BRIEF, extract, testing::Combine( testing::Values( BRIEF_IMAGES ), testing::Values( 16, 32, 64 ) ) )
+PERF_TEST_P( OCL_BRIEF, extract, testing::Combine(
+ testing::Values( string( "gpu/opticalflow/rubberwhale1.png" ),
+ string( "gpu/stereobm/aloe-L.png" )
+ ), testing::Values( 16, 32, 64 ), testing::Values( 250, 500, 1000, 2500, 3000 ) ) )
{
- 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 );
- ASSERT_FALSE( img.empty( ) );
+ const size_t numKp = std::tr1::get<2>(GetParam( ));
- if ( RUN_OCL_IMPL )
+ 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)
{
- oclMat d_img( img );
- oclMat d_keypoints;
- FAST_OCL fast( threshold );
- fast( d_img, oclMat( ), d_keypoints );
+ 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 )
+ {
+ 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;
+ }
BRIEF_OCL brief( bytes );
-
- OCL_TEST_CYCLE( )
+ oclMat imgCL( img ), keypointsCL(kpMat), mask;
+ while (next( ))
{
- oclMat d_descriptors;
- brief.compute( d_img, d_keypoints, d_descriptors );
+ startTimer( );
+ oclMat descriptorsCL;
+ brief.compute( imgCL, keypointsCL, mask, descriptorsCL );
+ cv::ocl::finish( );
+ stopTimer( );
}
-
- std::vector<KeyPoint> ocl_keypoints;
- fast.downloadKeypoints( d_keypoints, ocl_keypoints );
- SANITY_CHECK_KEYPOINTS( ocl_keypoints );
+ 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_KEYPOINTS( keypoints );
+ SANITY_CHECK_NOTHING( )
}
else
OCL_PERF_ELSE;
{
}
-void
-BRIEF_OCL::compute( const oclMat& image, oclMat& keypoints, oclMat& descriptors ) const
+void BRIEF_OCL::compute( const oclMat& image, const oclMat& keypoints, oclMat& mask, oclMat& descriptors ) const
{
- oclMat grayImage = image;
- if ( image.type( ) != CV_8U ) cvtColor( image, grayImage, COLOR_BGR2GRAY );
-
+ CV_Assert( image.type( ) == CV_8UC1 );
+ 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 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 ) );
-
- Context* ctx = Context::getContext( );
+ args.push_back( std::make_pair( sizeof (cl_mem), (void *) &mask.data ) );
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;
+}
//
//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)
#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 ))
- - read_imagei( sum, sampler, pt + (int2)( -HALF_KERNEL, HALF_KERNEL + 1 ))
- - read_imagei( sum, sampler, pt + (int2)( HALF_KERNEL + 1, -HALF_KERNEL ))
- + read_imagei( sum, sampler, pt + (int2)( -HALF_KERNEL, -HALF_KERNEL ))).x;
+ return ( read_imagei( sum, sampler, kpPos + pt + (int2)( HALF_KERNEL + 1, HALF_KERNEL + 1 ))
+ - read_imagei( sum, sampler, kpPos + pt + (int2)( -HALF_KERNEL, HALF_KERNEL + 1 ))
+ - read_imagei( sum, sampler, kpPos + pt + (int2)( HALF_KERNEL + 1, -HALF_KERNEL ))
+ + 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);
+ const int byte = get_local_id(0);
+ const int kpId = get_group_id(0);
+ 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, (int2)( tests[byte * 32 + (i * 4) + 2], tests[byte * 32 + (i * 4) + 3] ))
+ smoothedSum(sumImg, pt, (int2)( tests[byte * 32 + (i * 4) + 1], tests[byte * 32 + (i * 4) + 0] ))
+ < 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;
}
--- /dev/null
+/*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
\ No newline at end of file