Merge pull request #121 from mitjap/multi_gpu

add support for device selection and multiple GPUs

Author: griwodz

CHANGES.md

@@ -23,6 +23,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 - CMake: support for cuda 11 [PR](https://github.com/alicevision/popsift/pull/103)
 - Support for Cuda CC 7 cards (RTX 2080) [PR](https://github.com/alicevision/popsift/pull/67)
 - Support for Boost 1.70 [PR](https://github.com/alicevision/popsift/pull/65)
+- Support for device selection and multiple GPUs [PR](https://github.com/alicevision/popsift/pull/121)
 
 ### Fixed
 - CMake: fixes to allow building on Windows using vcpkg [PR](https://github.com/alicevision/popsift/pull/92)

src/popsift/gauss_filter.cu

@@ -18,7 +18,7 @@ namespace popsift {
 __device__ __constant__
 GaussInfo d_gauss;
 
-__align__(128) GaussInfo h_gauss;
+__align__(128) thread_local GaussInfo h_gauss;
 
 
 __global__

src/popsift/gauss_filter.h

@@ -105,7 +105,7 @@ struct GaussInfo
 };
 
 extern __device__ __constant__ GaussInfo d_gauss;
-extern                         GaussInfo h_gauss;
+extern thread_local            GaussInfo h_gauss;
 
 /* init_filter must be called early to initialize the Gauss tables.
  */

src/popsift/popsift.cpp

@@ -19,9 +19,13 @@
 
 using namespace std;
 
-PopSift::PopSift( const popsift::Config& config, popsift::Config::ProcessingMode mode, ImageMode imode )
+PopSift::PopSift( const popsift::Config& config, popsift::Config::ProcessingMode mode, ImageMode imode, int device )
     : _image_mode( imode )
+    , _device(device)
 {
+    cudaSetDevice(_device);
+    configure(config);
+
     if( imode == ByteImages )
     {
         _pipe._unused.push( new popsift::Image);
@@ -33,18 +37,19 @@ PopSift::PopSift( const popsift::Config& config, popsift::Config::ProcessingMode
         _pipe._unused.push( new popsift::ImageFloat );
     }
 
-    configure( config, true );
-
     _pipe._thread_stage1.reset( new std::thread( &PopSift::uploadImages, this ));
     if( mode == popsift::Config::ExtractingMode )
         _pipe._thread_stage2.reset( new std::thread( &PopSift::extractDownloadLoop, this ));
     else
         _pipe._thread_stage2.reset( new std::thread( &PopSift::matchPrepareLoop, this ));
 }
 
-PopSift::PopSift( ImageMode imode )
+PopSift::PopSift( ImageMode imode, int device )
     : _image_mode( imode )
+    , _device(device)
 {
+    cudaSetDevice(_device);
+
     if( imode == ByteImages )
     {
         _pipe._unused.push( new popsift::Image);
@@ -68,16 +73,20 @@ PopSift::~PopSift()
     }
 }
 
-bool PopSift::configure( const popsift::Config& config, bool force )
+bool PopSift::configure( const popsift::Config& config, bool /*force*/ )
 {
     if( _pipe._pyramid != nullptr ) {
         return false;
     }
 
     _config = config;
-
     _config.levels = max( 2, config.levels );
 
+    return true;
+}
+
+bool PopSift::applyConfiguration(bool force)
+{
     if( force || ( _config  != _shadow_config ) )
     {
         popsift::init_filter( _config,
@@ -131,6 +140,16 @@ bool PopSift::private_init( int w, int h )
     return true;
 }
 
+bool PopSift::private_uninit()
+{
+    Pipe& p = _pipe;
+
+    delete p._pyramid;
+    p._pyramid = nullptr;
+
+    return true;
+}
+
 void PopSift::uninit( )
 {
     if(!_isInit)
@@ -273,6 +292,8 @@ SiftJob* PopSift::enqueue( int          w,
 
 void PopSift::uploadImages( )
 {
+    cudaSetDevice(_device);
+
     SiftJob* job;
     while( ( job = _pipe._queue_stage1.pull() ) != nullptr ) {
         popsift::ImageBase* img = _pipe._unused.pull();
@@ -284,10 +305,15 @@ void PopSift::uploadImages( )
 
 void PopSift::extractDownloadLoop( )
 {
+    cudaSetDevice(_device);
+    applyConfiguration(true);
+
     Pipe& p = _pipe;
 
     SiftJob* job;
     while( ( job = p._queue_stage2.pull() ) != nullptr ) {
+        applyConfiguration();
+
         popsift::ImageBase* img = job->getImg();
 
         private_init( img->getWidth(), img->getHeight() );
@@ -313,14 +339,21 @@ void PopSift::extractDownloadLoop( )
 
         job->setFeatures( features );
     }
+
+    private_uninit();
 }
 
 void PopSift::matchPrepareLoop( )
 {
+    cudaSetDevice(_device);
+    applyConfiguration(true);
+
     Pipe& p = _pipe;
 
     SiftJob* job;
     while( ( job = p._queue_stage2.pull() ) != nullptr ) {
+        applyConfiguration();
+
         popsift::ImageBase* img = job->getImg();
 
         private_init( img->getWidth(), img->getHeight() );
@@ -336,6 +369,8 @@ void PopSift::matchPrepareLoop( )
 
         job->setFeatures( features );
     }
+
+    private_uninit();
 }
 
 SiftJob::SiftJob( int w, int h, const unsigned char* imageData )
@@ -445,8 +480,4 @@ void PopSift::Pipe::uninit()
         popsift::ImageBase* img = _unused.pull();
         delete img;
     }
-
-    delete _pyramid;
-    _pyramid    = nullptr;
-
 }

src/popsift/popsift.h

@@ -150,7 +150,7 @@ class PopSift
      * @brief We support more than 1 streams, but we support only one sigma and one
      * level parameters.
      */
-    explicit PopSift( ImageMode imode = ByteImages );
+    explicit PopSift( ImageMode imode = ByteImages, int device = 0 );
 
     /**
      * @brief
@@ -160,7 +160,7 @@ class PopSift
      */
     explicit PopSift(const popsift::Config& config,
                      popsift::Config::ProcessingMode mode = popsift::Config::ExtractingMode,
-                     ImageMode imode = ByteImages);
+                     ImageMode imode = ByteImages, int device = 0);
 
     /**
      * @brief Release all the resources.
@@ -273,7 +273,10 @@ class PopSift
     }
 
 private:
+    bool applyConfiguration( bool force = false );
+
     bool private_init( int w, int h );
+    bool private_uninit( );
     void private_apply_scale_factor( int& w, int& h );
     void uploadImages( );
 
@@ -299,6 +302,7 @@ class PopSift
     int             _last_init_w{}; /* to support deprecated interface */
     int             _last_init_h{}; /* to support deprecated interface */
     ImageMode       _image_mode;
+    int             _device;
 
     /// whether the object is initialized
     bool            _isInit{true};

src/popsift/sift_constants.cu

@@ -16,7 +16,7 @@ using namespace std;
 
 namespace popsift {
 
-ConstInfo                         h_consts;
+thread_local            ConstInfo h_consts;
 __device__ __constant__ ConstInfo d_consts;
 
 void init_constants( float sigma0, int levels, float threshold, float edge_limit, int max_extrema, int normalization_multiplier )

src/popsift/sift_constants.h

@@ -68,7 +68,7 @@ struct ConstInfo
     float desc_tile[16];
 };
 
-extern                         ConstInfo h_consts;
+extern thread_local            ConstInfo h_consts;
 extern __device__ __constant__ ConstInfo d_consts;
 
 

src/popsift/sift_pyramid.cu

@@ -38,18 +38,15 @@ using namespace std;
 
 namespace popsift {
 
-__device__
-ExtremaCounters dct;
-ExtremaCounters hct;
+__device__ ExtremaCounters   dct;
+thread_local ExtremaCounters hct;
 
-__device__
-ExtremaBuffers  dbuf;
-ExtremaBuffers  dbuf_shadow; // just for managing memories
-ExtremaBuffers  hbuf;
+__device__ ExtremaBuffers   dbuf;
+thread_local ExtremaBuffers dbuf_shadow; // just for managing memories
+thread_local ExtremaBuffers hbuf;
 
-__device__
-DevBuffers      dobuf;
-DevBuffers      dobuf_shadow; // just for managing memories
+__device__ DevBuffers       dobuf;
+thread_local DevBuffers     dobuf_shadow; // just for managing memories
 
 __global__
     void py_print_corner_float(float* img, uint32_t pitch, uint32_t height, uint32_t level)
@@ -215,6 +212,7 @@ Pyramid::~Pyramid()
 {
     cudaStreamDestroy( _download_stream );
 
+    cudaFree(     _d_extrema_num_blocks );
     cudaFree(     dobuf_shadow.i_ext_dat[0] );
     cudaFree(     dobuf_shadow.i_ext_off[0] );
     cudaFree(     dobuf_shadow.features );

src/popsift/sift_pyramid.h

@@ -50,13 +50,13 @@ struct DevBuffers
     Feature*         features;
 };
 
-extern            ExtremaCounters hct;
-extern __device__ ExtremaCounters dct;
-extern            ExtremaBuffers  hbuf;
-extern __device__ ExtremaBuffers  dbuf;
-extern            ExtremaBuffers  dbuf_shadow; // just for managing memories
-extern __device__ DevBuffers      dobuf;
-extern            DevBuffers      dobuf_shadow; // just for managing memories
+extern thread_local ExtremaCounters hct;
+extern __device__   ExtremaCounters dct;
+extern thread_local ExtremaBuffers  hbuf;
+extern __device__   ExtremaBuffers  dbuf;
+extern thread_local ExtremaBuffers  dbuf_shadow; // just for managing memories
+extern __device__   DevBuffers      dobuf;
+extern thread_local DevBuffers      dobuf_shadow; // just for managing memories
 
 class Pyramid
 {