mirror of https://github.com/opencv/opencv.git
Merge pull request #2027 from AD-530:brief_cl
commit
06acf7090c
8 changed files with 596 additions and 32 deletions
@ -0,0 +1,114 @@ |
||||
/*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) 2013, OpenCV Foundation, all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||
// this list of conditions and the following disclaimer in the documentation
|
||||
// and/or other materials provided with the distribution.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// This software is provided by the copyright holders and contributors "as is" and
|
||||
// any express or implied warranties, including, but not limited to, the implied
|
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||
// indirect, incidental, special, exemplary, or consequential damages
|
||||
// (including, but not limited to, procurement of substitute goods or services;
|
||||
// loss of use, data, or profits; or business interruption) however caused
|
||||
// and on any theory of liability, whether in contract, strict liability,
|
||||
// or tort (including negligence or otherwise) arising in any way out of
|
||||
// the use of this software, even if advised of the possibility of such damage.
|
||||
//
|
||||
// Authors:
|
||||
// * Matthias Bady, aegirxx ==> gmail.com
|
||||
//
|
||||
//M*/
|
||||
|
||||
#include "perf_precomp.hpp" |
||||
|
||||
using namespace std; |
||||
using namespace cv; |
||||
using namespace ocl; |
||||
using namespace perf; |
||||
|
||||
///////////// BRIEF ////////////////////////
|
||||
typedef TestBaseWithParam<std::tr1::tuple<std::string, int, size_t> > OCL_BRIEF; |
||||
|
||||
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 std::string filename = std::tr1::get<0>(GetParam( )); |
||||
const int bytes = std::tr1::get<1>(GetParam( )); |
||||
const size_t numKp = std::tr1::get<2>(GetParam( )); |
||||
|
||||
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 ) |
||||
{ |
||||
Mat kpMat( 2, int( keypoints.size() ), CV_32FC1 ); |
||||
for ( size_t i = 0; i < keypoints.size( ); ++i ) |
||||
{ |
||||
kpMat.col( int( i ) ).row( 0 ) = keypoints[i].pt.x; |
||||
kpMat.col( int( i ) ).row( 1 ) = keypoints[i].pt.y; |
||||
} |
||||
BRIEF_OCL brief( bytes ); |
||||
oclMat imgCL( img ), keypointsCL(kpMat), mask; |
||||
while (next( )) |
||||
{ |
||||
startTimer( ); |
||||
oclMat descriptorsCL; |
||||
brief.compute( imgCL, keypointsCL, mask, descriptorsCL ); |
||||
cv::ocl::finish( ); |
||||
stopTimer( ); |
||||
} |
||||
SANITY_CHECK_NOTHING( ) |
||||
} |
||||
else if ( RUN_PLAIN_IMPL ) |
||||
{ |
||||
BriefDescriptorExtractor brief( bytes ); |
||||
|
||||
while (next( )) |
||||
{ |
||||
startTimer( ); |
||||
Mat descriptors; |
||||
brief.compute( img, keypoints, descriptors ); |
||||
stopTimer( ); |
||||
} |
||||
SANITY_CHECK_NOTHING( ) |
||||
} |
||||
else |
||||
OCL_PERF_ELSE; |
||||
} |
@ -0,0 +1,91 @@ |
||||
/*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 "precomp.hpp" |
||||
#include "opencl_kernels.hpp" |
||||
|
||||
using namespace cv; |
||||
using namespace cv::ocl; |
||||
|
||||
BRIEF_OCL::BRIEF_OCL( int _bytes ) : bytes( _bytes ) |
||||
{ |
||||
} |
||||
|
||||
void BRIEF_OCL::compute( const oclMat& image, const oclMat& keypoints, oclMat& mask, oclMat& descriptors ) const |
||||
{ |
||||
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( image, sum, CV_32S ); |
||||
cl_mem sumTexture = bindTexture( sum ); |
||||
std::stringstream build_opt; |
||||
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 ) ); |
||||
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; |
||||
} |
@ -0,0 +1,173 @@ |
||||
/*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 |
||||
// |
||||
// 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*/ |
||||
|
||||
#ifndef BYTES |
||||
#define BYTES 16 |
||||
#endif |
||||
|
||||
#ifndef KERNEL_SIZE |
||||
#define KERNEL_SIZE 9 |
||||
#endif |
||||
|
||||
#ifndef BORDER |
||||
#define BORDER 0 |
||||
#endif |
||||
|
||||
#define HALF_KERNEL (KERNEL_SIZE/2) |
||||
|
||||
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; |
||||
|
||||
__constant char tests[32 * BYTES] = |
||||
{ |
||||
-2,-1,7,-1,-14,-1,-3,3,1,-2,11,2,1,6,-10,-7,13,2,-1,0,-14,5,5,-3,-2,8,2,4,-11,8,-15,5, |
||||
-6,-23,8,-9,-12,6,-10,8,-3,-1,8,1,3,6,5,6,-7,-6,5,-5,22,-2,-11,-8,14,7,8,5,-1,14,-5,-14, |
||||
-14,9,2,0,7,-3,22,6,-6,6,-8,-5,-5,9,7,-1,-3,-7,-10,-18,4,-5,0,11,2,3,9,10,-10,3,4,9, |
||||
0,12,-3,19,1,15,-11,-5,14,-1,7,8,7,-23,-5,5,0,-6,-10,17,13,-4,-3,-4,-12,1,-12,2,0,8,3,22, |
||||
-13,13,3,-1,-16,17,6,10,7,15,-5,0,2,-12,19,-2,3,-6,-4,-15,8,3,0,14,4,-11,5,5,11,-7,7,1, |
||||
6,12,21,3,-3,2,14,1,5,1,-5,11,3,-17,-6,2,6,8,5,-10,-14,-2,0,4,5,-7,-6,5,10,4,4,-7, |
||||
22,0,7,-18,-1,-3,0,18,-4,22,-5,3,1,-7,2,-3,19,-20,17,-2,3,-10,-8,24,-5,-14,7,5,-2,12,-4,-15, |
||||
4,12,0,-19,20,13,3,5,-8,-12,5,0,-5,6,-7,-11,6,-11,-3,-22,15,4,10,1,-7,-4,15,-6,5,10,0,24, |
||||
3,6,22,-2,-13,14,4,-4,-13,8,-18,-22,-1,-1,-7,3,-19,-12,4,3,8,10,13,-2,-6,-1,-6,-5,2,-21,-3,2, |
||||
4,-7,0,16,-6,-5,-12,-1,1,-1,9,18,-7,10,-11,6,4,3,19,-7,-18,5,-4,5,4,0,-20,4,7,-11,18,12, |
||||
-20,17,-18,7,2,15,19,-11,-18,6,-7,3,-4,1,-14,13,17,3,2,-8,-7,2,1,6,17,-9,-2,8,-8,-6,-1,12, |
||||
-2,4,-1,6,-2,7,6,8,-8,-1,-7,-9,8,-9,15,0,0,22,-4,-15,-14,-1,3,-2,-7,-4,17,-7,-8,-2,9,-4, |
||||
5,-7,7,7,-5,13,-8,11,11,-4,0,8,5,-11,-9,-6,2,-6,3,-20,-6,2,6,10,-6,-6,-15,7,-6,-3,2,1, |
||||
11,0,-3,2,7,-12,14,5,0,-7,-1,-1,-16,0,6,8,22,11,0,-3,19,0,5,-17,-23,-14,-13,-19,-8,10,-11,-2, |
||||
-11,6,-10,13,1,-7,14,0,-12,1,-5,-5,4,7,8,-1,-1,-5,15,2,-3,-1,7,-10,3,-6,10,-18,-7,-13,-13,10, |
||||
1,-1,13,-10,-19,14,8,-14,-4,-13,7,1,1,-2,12,-7,3,-5,1,-5,-2,-2,8,-10,2,14,8,7,3,9,8,2 |
||||
#if BYTES > 16 |
||||
,-9,1,-18,0,4,0,1,12,0,9,-14,-10,-13,-9,-2,6,1,5,10,10,-3,-6,-16,-5,11,6,-5,0,-23,10,1,2, |
||||
13,-5,-3,9,-4,-1,-13,-5,10,13,-11,8,19,20,-9,2,4,-8,0,-9,-14,10,15,19,-14,-12,-10,-3,-23,-3,17,-2, |
||||
-3,-11,6,-14,19,-2,-4,2,-5,5,3,-13,2,-2,-5,4,17,4,17,-11,-7,-2,1,23,8,13,1,-16,-13,-5,1,-17, |
||||
4,6,-8,-3,-5,-9,-2,-10,-9,0,-7,-2,5,0,5,2,-4,-16,6,3,2,-15,-2,12,4,-1,6,2,1,1,-2,-8, |
||||
-2,12,-5,-2,-8,8,-9,9,2,-10,3,1,-4,10,-9,4,6,12,2,5,-3,-8,0,5,-13,1,-7,2,-1,-10,7,-18, |
||||
-1,8,-9,-10,-23,-1,6,2,-5,-3,3,2,0,11,-4,-7,15,2,-10,-3,-20,-8,-13,3,-19,-12,5,-11,-17,-13,-3,2, |
||||
7,4,-12,0,5,-1,-14,-6,-4,11,0,-4,3,10,7,-3,13,21,-11,6,-12,24,-7,-4,4,16,3,-14,-3,5,-7,-12, |
||||
0,-4,7,-5,-17,-9,13,-7,22,-6,-11,5,2,-8,23,-11,7,-10,-1,14,-3,-10,8,3,-13,1,-6,0,-7,-21,6,-14, |
||||
18,19,-4,-6,10,7,-1,-4,-1,21,1,-5,-10,6,-11,-2,18,-3,-1,7,-3,-9,-5,10,-13,14,17,-3,11,-19,-1,-18, |
||||
8,-2,-18,-23,0,-5,-2,-9,-4,-11,2,-8,14,6,-3,-6,-3,0,-15,0,-9,4,-15,-9,-1,11,3,11,-10,-16,-7,7, |
||||
-2,-10,-10,-2,-5,-3,5,-23,13,-8,-15,-11,-15,11,6,-6,-16,-3,-2,2,6,12,-16,24,-10,0,8,11,-7,7,-19,-7, |
||||
5,16,9,-3,9,7,-7,-16,3,2,-10,9,21,1,8,7,7,0,1,17,-8,12,9,6,11,-7,-8,-6,19,0,9,3, |
||||
1,-7,-5,-11,0,8,-2,14,12,-2,-15,-6,4,12,0,-21,17,-4,-6,-7,-10,-9,-14,-7,-15,-10,-15,-14,-7,-5,5,-12, |
||||
-4,0,15,-4,5,2,-6,-23,-4,-21,-6,4,-10,5,-15,6,4,-3,-1,5,-4,19,-23,-4,-4,17,13,-11,1,12,4,-14, |
||||
-11,-6,-20,10,4,5,3,20,-8,-20,3,1,-19,9,9,-3,18,15,11,-4,12,16,8,7,-14,-8,-3,9,-6,0,2,-4, |
||||
1,-10,-1,2,8,-7,-6,18,9,12,-7,-23,8,-6,5,2,-9,6,-12,-7,-1,-2,-7,2,9,9,7,15,6,2,-6,6 |
||||
#endif |
||||
#if BYTES > 32 |
||||
,16,12,0,19,4,3,6,0,-2,-1,2,17,8,1,3,1,-12,-1,-11,0,-11,2,7,9,-1,3,-19,4,-1,-11,-1,3, |
||||
1,-10,-10,-4,-2,3,6,11,3,7,-9,-8,24,-14,-2,-10,-3,-3,-18,-6,-13,-10,-7,-1,2,-7,9,-6,2,-4,6,-13, |
||||
4,-4,-2,3,-4,2,9,13,-11,5,-6,-11,4,-2,11,-9,-19,0,-23,-5,-5,-7,-3,-6,-6,-4,12,14,12,-11,-8,-16, |
||||
-21,15,-12,6,-2,-1,-8,16,6,-1,-8,-2,1,-1,-9,8,3,-4,-2,-2,-7,0,4,-8,11,-11,-12,2,2,3,11,7, |
||||
-7,-4,-9,-6,3,-7,-5,0,3,-7,-10,-5,-3,-1,8,-10,0,8,5,1,9,0,1,16,8,4,-11,-3,-15,9,8,17, |
||||
0,2,-9,17,-6,-11,-10,-3,1,1,15,-8,-12,-13,-2,4,-6,4,-6,-10,5,-7,7,-5,10,6,8,9,-5,7,-18,-3, |
||||
-6,3,5,4,-10,-13,-5,-3,-11,2,-16,0,7,-21,-5,-13,-14,-14,-4,-4,4,9,7,-3,4,11,10,-4,6,17,9,17, |
||||
-10,8,0,-11,-6,-16,-6,8,-13,5,10,-5,3,2,12,16,13,-8,0,-6,10,0,4,-11,8,5,10,-2,11,-7,-13,3, |
||||
2,4,-7,-3,-14,-2,-11,16,11,-6,7,6,-3,15,8,-10,-3,8,12,-12,-13,6,-14,7,-11,-5,-8,-6,7,-6,6,3, |
||||
-4,10,5,1,9,16,10,13,-17,10,2,8,-5,1,4,-4,-14,8,-5,2,4,-9,-6,-3,3,-7,-10,0,-2,-8,-10,4, |
||||
-8,5,-9,24,2,-8,8,-9,-4,17,-5,2,14,0,-9,9,11,15,-6,5,-8,1,-3,4,9,-21,10,2,2,-1,4,11, |
||||
24,3,2,-2,-8,17,-14,-10,6,5,-13,7,11,10,0,-1,4,6,-10,6,-12,-2,5,6,3,-1,8,-15,1,-4,-7,11, |
||||
1,11,5,0,6,-12,10,1,-3,-2,-1,4,-2,-11,-1,12,7,-8,-20,-18,2,0,-9,2,-13,-1,-16,2,3,-1,-5,-17, |
||||
15,8,3,-14,-13,-12,6,15,2,-8,2,6,6,22,-3,-23,-2,-7,-6,0,13,-10,-6,6,6,7,-10,12,-6,7,-2,11, |
||||
0,-22,-2,-17,-4,-1,-11,-14,-2,-8,7,12,12,-5,7,-13,2,-2,-7,6,0,8,-3,23,6,12,13,-11,-21,-10,10,8, |
||||
-3,0,7,15,7,-6,-5,-12,-21,-10,12,-11,-5,-11,8,-11,5,0,-11,-1,8,-9,7,-1,11,-23,21,-5,0,-5,-8,6, |
||||
-6,8,8,12,-7,5,3,-2,-5,-20,-12,9,-6,12,-11,3,4,5,13,11,2,12,13,-12,-4,-13,4,7,0,15,-3,-16, |
||||
-3,2,-2,14,4,-14,16,-11,-13,3,23,10,9,-19,2,5,5,3,14,-7,19,-13,-11,15,14,0,-2,-5,11,-4,0,-6, |
||||
-2,5,-13,-8,-11,-15,-7,-17,1,3,-10,-8,-13,-10,7,-12,0,-13,23,-6,2,-17,-7,-3,1,3,4,-10,13,4,14,-6, |
||||
-19,-2,-1,5,9,-8,10,-5,7,-1,5,7,9,-10,19,0,7,5,-4,-7,-11,1,-1,-11,2,-1,-4,11,-1,7,2,-2, |
||||
1,-20,-9,-6,-4,-18,8,-18,-16,-2,7,-6,-3,-6,-1,-4,0,-16,24,-5,-4,-2,-1,9,-8,2,-6,15,11,4,0,-3, |
||||
7,6,2,-10,-7,-9,12,-6,24,15,-8,-1,15,-9,-3,-15,17,-5,11,-10,-2,13,-15,4,-2,-1,4,-23,-16,3,-7,-14, |
||||
-3,-5,-10,-9,-5,3,-2,-1,-1,4,1,8,12,9,9,-14,-9,17,-3,0,5,4,13,-6,-1,-8,19,10,8,-5,-15,2, |
||||
-12,-9,-4,-5,12,0,24,4,8,-2,14,4,8,-4,-7,16,5,-1,-8,-4,-2,18,-5,17,8,-2,-9,-2,3,-7,1,-6, |
||||
-5,-22,-5,-2,-8,-10,14,1,-3,-13,3,9,-4,-1,-1,0,-7,-21,12,-19,-8,8,24,8,12,-6,-2,3,-5,-11,-22,-4, |
||||
-3,5,-4,4,-16,24,7,-9,-10,23,-9,18,1,12,17,21,24,-6,-3,-11,-7,17,1,-6,4,4,2,-7,14,6,-12,3, |
||||
-6,0,-16,13,-10,5,7,12,5,2,6,-3,7,0,-23,1,15,-5,1,14,-3,-1,6,6,6,-9,-9,12,4,-2,-4,7, |
||||
-4,-5,4,4,-13,0,6,-10,2,-12,-6,-3,16,0,-3,3,5,-14,6,11,5,11,0,-13,7,5,-1,-5,12,4,6,10, |
||||
-10,4,-1,-11,4,10,-14,5,11,-14,-13,0,2,8,12,24,-1,3,-1,2,9,-14,-23,3,-8,-6,0,9,-15,14,10,-10, |
||||
-10,-6,-7,-5,11,5,-3,-15,1,0,1,8,-11,-6,-4,-18,9,0,22,-4,-5,-1,-9,4,-20,2,1,6,1,2,-9,-12, |
||||
5,15,4,-6,19,4,4,11,17,-4,-8,-1,-8,-12,7,-3,11,9,8,1,9,22,-15,15,-7,-7,1,-23,-5,13,-8,2, |
||||
3,-5,11,-11,3,-18,14,-5,-20,7,-10,-23,-2,-5,6,0,-17,-13,-3,2,-6,-1,14,-2,-12,-16,15,6,-12,-2,3,-19 |
||||
#endif |
||||
}; |
||||
|
||||
inline int smoothedSum(__read_only image2d_t sum, const int2 kpPos, const int2 pt) |
||||
{ |
||||
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, |
||||
__global uchar* mask) |
||||
{ |
||||
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, 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; |
||||
} |
@ -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, int( keypoints.size() ), CV_32FC1 ); |
||||
for ( size_t i = 0; i < keypoints.size( ); ++i ) |
||||
{ |
||||
kpMat.col( i ).row( 0 ) = int( keypoints[i].pt.x ); |
||||
kpMat.col( i ).row( 1 ) = int( 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 ) |
||||
{ |
||||
int 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 |
Loading…
Reference in new issue