diff --git a/src/popsift/common/debug_macros.cu b/src/popsift/common/debug_macros.cu index cf4cd735..c9155248 100755 --- a/src/popsift/common/debug_macros.cu +++ b/src/popsift/common/debug_macros.cu @@ -20,12 +20,7 @@ void pop_sync_check_last_error( const char* file, size_t line ) void pop_check_last_error( const char* file, size_t line ) { cudaError_t err = cudaGetLastError( ); - if( err != cudaSuccess ) { - std::cerr << __FILE__ << ":" << __LINE__ << std::endl - << " called from " << file << ":" << line << std::endl - << " cudaGetLastError failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaGetLastError failed: "); } namespace popsift { namespace cuda { @@ -34,11 +29,7 @@ void malloc_dev( void** ptr, int sz, { cudaError_t err; err = cudaMalloc( ptr, sz ); - if( err != cudaSuccess ) { - std::cerr << file << ":" << line << std::endl - << " cudaMalloc failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaMalloc failed: "); #ifdef DEBUG_INIT_DEVICE_ALLOCATIONS popsift::cuda::memset_sync( *ptr, 0, sz, file, line ); #endif // NDEBUG @@ -51,11 +42,7 @@ void malloc_hst( void** ptr, int sz, { cudaError_t err; err = cudaMallocHost( ptr, sz ); - if( err != cudaSuccess ) { - std::cerr << file << ":" << line << std::endl - << " cudaMallocHost failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaMallocHost failed: "); #ifdef DEBUG_INIT_DEVICE_ALLOCATIONS memset( *ptr, 0, sz ); #endif // NDEBUG @@ -74,16 +61,13 @@ void memcpy_async( void* dst, const void* src, size_t sz, cudaError_t err; err = cudaMemcpyAsync( dst, src, sz, type, stream ); if( err != cudaSuccess ) { - cerr << file << ":" << line << endl - << " " << "Failed to copy " - << (type==cudaMemcpyHostToDevice?"host-to-device":"device-to-host") - << ": "; - cerr << cudaGetErrorString(err) << endl; - cerr << " src ptr=" << hex << (size_t)src << dec << endl - << " dst ptr=" << hex << (size_t)dst << dec << endl; - exit( -__LINE__ ); + std::stringstream ss; + ss << "Failed to copy " << (type == cudaMemcpyHostToDevice ? "host-to-device" : "device-to-host") << ": "; + ss << cudaGetErrorString(err) << endl; + ss << " src ptr=" << hex << (size_t)src << dec << endl + << " dst ptr=" << hex << (size_t)dst << dec << endl; + POP_FATAL(ss.str()); } - POP_CUDA_FATAL_TEST( err, "Failed to copy host-to-device: " ); } void memcpy_sync( void* dst, const void* src, size_t sz, cudaMemcpyKind type, const char* file, size_t line ) @@ -95,37 +79,27 @@ void memcpy_sync( void* dst, const void* src, size_t sz, cudaMemcpyKind type, co cudaError_t err; err = cudaMemcpy( dst, src, sz, type ); if( err != cudaSuccess ) { - cerr << " " << "Failed to copy " - << (type==cudaMemcpyHostToDevice?"host-to-device":"device-to-host") - << ": "; - cerr << cudaGetErrorString(err) << endl; - cerr << " src ptr=" << hex << (size_t)src << dec << endl - << " dst ptr=" << hex << (size_t)dst << dec << endl; - exit( -__LINE__ ); + std::stringstream ss; + ss << "Failed to copy " << (type == cudaMemcpyHostToDevice ? "host-to-device" : "device-to-host") << ": "; + ss << cudaGetErrorString(err) << endl; + ss << " src ptr=" << hex << (size_t)src << dec << endl + << " dst ptr=" << hex << (size_t)dst << dec << endl; + POP_FATAL(ss.str()) } - POP_CUDA_FATAL_TEST( err, "Failed to copy host-to-device: " ); } void memset_async( void* ptr, int value, size_t bytes, cudaStream_t stream, const char* file, size_t line ) { cudaError_t err; err = cudaMemsetAsync( ptr, value, bytes, stream ); - if( err != cudaSuccess ) { - std::cerr << file << ":" << line << std::endl - << " cudaMemsetAsync failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaMemsetAsync failed: "); } void memset_sync( void* ptr, int value, size_t bytes, const char* file, size_t line ) { cudaError_t err; err = cudaMemset( ptr, value, bytes ); - if( err != cudaSuccess ) { - std::cerr << file << ":" << line << std::endl - << " cudaMemset failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaMemset failed: "); } } } @@ -135,68 +109,44 @@ cudaStream_t stream_create( const char* file, size_t line ) cudaStream_t stream; cudaError_t err; err = cudaStreamCreate( &stream ); - if( err != cudaSuccess ) { - std::cerr << file << ":" << line << std::endl - << " cudaStreamCreate failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaStreamCreate failed: "); return stream; } void stream_destroy( cudaStream_t s, const char* file, size_t line ) { cudaError_t err; err = cudaStreamDestroy( s ); - if( err != cudaSuccess ) { - std::cerr << file << ":" << line << std::endl - << " cudaStreamDestroy failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaStreamDestroy failed: "); } cudaEvent_t event_create( const char* file, size_t line ) { cudaEvent_t ev; cudaError_t err; err = cudaEventCreate( &ev ); - if( err != cudaSuccess ) { - std::cerr << file << ":" << line << std::endl - << " cudaEventCreate failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaEventCreate failed: "); return ev; } void event_destroy( cudaEvent_t ev, const char* file, size_t line ) { cudaError_t err; err = cudaEventDestroy( ev ); - if( err != cudaSuccess ) { - std::cerr << file << ":" << line << std::endl - << " cudaEventDestroy failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaEventDestroy failed: "); } void event_record( cudaEvent_t ev, cudaStream_t s, const char* file, size_t line ) { cudaError_t err; err = cudaEventRecord( ev, s ); - if( err != cudaSuccess ) { - std::cerr << file << ":" << line << std::endl - << " cudaEventRecord failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaEventRecord failed: "); } void event_wait( cudaEvent_t ev, cudaStream_t s, const char* file, size_t line ) { cudaError_t err; err = cudaStreamWaitEvent( s, ev, 0 ); - if( err != cudaSuccess ) { - std::cerr << file << ":" << line << std::endl - << " cudaStreamWaitEvent failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaStreamWaitEvent failed: "); } float event_diff( cudaEvent_t from, cudaEvent_t to ) -{ +{ float ms; cudaEventElapsedTime( &ms, from, to ); return ms; diff --git a/src/popsift/common/debug_macros.h b/src/popsift/common/debug_macros.h index dada7d3b..c8f31651 100755 --- a/src/popsift/common/debug_macros.h +++ b/src/popsift/common/debug_macros.h @@ -13,6 +13,7 @@ #include #include #include +#include #include // synchronize device and check for an error @@ -117,14 +118,18 @@ class BriefDuration }; }; -#define POP_FATAL(s) { \ - std::cerr << __FILE__ << ":" << __LINE__ << std::endl << " " << s << std::endl; \ - exit( -__LINE__ ); \ +#define POP_FATAL(s) \ + { \ + std::stringstream ss; \ + ss << __FILE__ << ":" << __LINE__ << std::endl << " " << s; \ + throw std::runtime_error{ss.str()}; \ } -#define POP_FATAL_FL(s,file,line) { \ - std::cerr << file << ":" << line << std::endl << " " << s << std::endl; \ - exit( -__LINE__ ); \ +#define POP_FATAL_FL(s, file, line) \ + { \ + std::stringstream ss; \ + ss << file << ":" << line << std::endl << " " << s << std::endl; \ + throw std::runtime_error{ss.str()}; \ } #define POP_CHECK_NON_NULL(ptr,s) if( ptr == 0 ) { POP_FATAL_FL(s,__FILE__,__LINE__); } @@ -147,10 +152,12 @@ class BriefDuration std::cerr << __FILE__ << ":" << __LINE__ << std::endl; \ std::cerr << " WARNING: " << s << cudaGetErrorString(err) << std::endl; \ } -#define POP_CUDA_FATAL(err,s) { \ - std::cerr << __FILE__ << ":" << __LINE__ << std::endl; \ - std::cerr << " " << s << cudaGetErrorString(err) << std::endl; \ - exit( -__LINE__ ); \ +#define POP_CUDA_FATAL(err,s) \ + { \ + std::stringstream ss; \ + ss << __FILE__ << ":" << __LINE__ << std::endl; \ + ss << " " << s << cudaGetErrorString(err) << std::endl; \ + throw std::runtime_error{ss.str()}; \ } #define POP_CUDA_FATAL_TEST(err,s) if( err != cudaSuccess ) { POP_CUDA_FATAL(err,s); } diff --git a/src/popsift/common/plane_2d.cu b/src/popsift/common/plane_2d.cu index 773dac24..89ba7d34 100644 --- a/src/popsift/common/plane_2d.cu +++ b/src/popsift/common/plane_2d.cu @@ -15,6 +15,7 @@ #include #include #include +#include #ifndef _WIN32 #include #else @@ -65,11 +66,11 @@ void* PlaneBase::allocHost2D( int w, int h, int elemSize, PlaneMapMode m ) #else const char *buf = strerror(errno); #endif - cerr << __FILE__ << ":" << __LINE__ << endl - << " Failed to allocate " << sz << " bytes of unaligned host memory." << endl - << " Cause: " << buf << endl; - exit( -1 ); - } else if( m == PageAligned ) { + stringstream ss; + ss << "Failed to allocate " << sz << " bytes of unaligned host memory." << endl + << "Cause: " << buf; + POP_FATAL(ss.str()); + } else if(m == PageAligned) { void* ptr = memalign(getPageSize(), sz); if(ptr) return ptr; @@ -93,9 +94,7 @@ void* PlaneBase::allocHost2D( int w, int h, int elemSize, PlaneMapMode m ) POP_CUDA_FATAL_TEST( err, "Failed to allocate aligned and pinned host memory: " ); return ptr; } else { - cerr << __FILE__ << ":" << __LINE__ << endl - << " Alignment not correctly specified in host plane allocation" << endl; - exit( -1 ); + POP_FATAL("Alignment not correctly specified in host plane allocation"); } } diff --git a/src/popsift/common/plane_2d.h b/src/popsift/common/plane_2d.h index 86d26f5c..2dad48cd 100644 --- a/src/popsift/common/plane_2d.h +++ b/src/popsift/common/plane_2d.h @@ -16,6 +16,8 @@ #include #include +#include "debug_macros.h" + #define PLANE2D_CUDA_OP_DEBUG #ifndef NDEBUG @@ -407,14 +409,16 @@ __host__ inline void Plane2D::memcpyToDevice( Plane2D& devPlane, cudaStream_t stream ) { if( devPlane._cols != this->_cols ) { - std::cerr << __FILE__ << ":" << __LINE__ << std::endl - << " Error: source columns (" << this->_cols << ") and dest columns (" << devPlane._cols << ") must be identical" << std::endl; - exit( -1 ); + std::stringstream ss; + ss << "Error: source columns (" << this->_cols << ") and dest columns (" << devPlane._cols + << ") must be identical"; + POP_FATAL(ss.str()); } if( devPlane._rows != this->_rows ) { - std::cerr << __FILE__ << ":" << __LINE__ << std::endl - << " Error: source rows (" << this->_rows << ") and dest rows (" << devPlane._rows << ") must be identical" << std::endl; - exit( -1 ); + std::stringstream ss; + ss << "Error: source rows (" << this->_rows << ") and dest rows (" << devPlane._rows + << ") must be identical"; + POP_FATAL(ss.str()); } PitchPlane2D::memcpyToDevice( devPlane, this->_cols, this->_rows, stream ); } diff --git a/src/popsift/features.cu b/src/popsift/features.cu index 023279ff..5aa706a1 100755 --- a/src/popsift/features.cu +++ b/src/popsift/features.cu @@ -16,6 +16,7 @@ #include #include #include +#include using namespace std; @@ -61,19 +62,21 @@ void FeaturesHost::reset( int num_ext, int num_ori ) _ext = (Feature*)memalign( getPageSize(), num_ext * sizeof(Feature) ); if( _ext == nullptr ) { - cerr << __FILE__ << ":" << __LINE__ << " Runtime error:" << endl - << " Failed to (re)allocate memory for downloading " << num_ext << " features" << endl; - if( errno == EINVAL ) cerr << " Alignment is not a power of two." << endl; - if( errno == ENOMEM ) cerr << " Not enough memory." << endl; - exit( -1 ); + std::stringstream ss; + ss << "Runtime error:" << endl + << " Failed to (re)allocate memory for downloading " << num_ext << " features" << endl; + if(errno == EINVAL) ss << " Alignment is not a power of two."; + if(errno == ENOMEM) ss << " Not enough memory."; + POP_FATAL(ss.str()); } _ori = (Descriptor*)memalign( getPageSize(), num_ori * sizeof(Descriptor) ); - if( _ori == nullptr ) { - cerr << __FILE__ << ":" << __LINE__ << " Runtime error:" << endl - << " Failed to (re)allocate memory for downloading " << num_ori << " descriptors" << endl; - if( errno == EINVAL ) cerr << " Alignment is not a power of two." << endl; - if( errno == ENOMEM ) cerr << " Not enough memory." << endl; - exit( -1 ); + if(_ori == nullptr) { + std::stringstream ss; + ss << "Runtime error:" << endl + << " Failed to (re)allocate memory for downloading " << num_ori << " descriptors" << endl; + if(errno == EINVAL) ss << " Alignment is not a power of two."; + if(errno == ENOMEM) ss << " Not enough memory."; + POP_FATAL(ss.str()); } setFeatureCount( num_ext ); diff --git a/src/popsift/gauss_filter.cu b/src/popsift/gauss_filter.cu index 537c843e..7c425f7f 100755 --- a/src/popsift/gauss_filter.cu +++ b/src/popsift/gauss_filter.cu @@ -130,17 +130,17 @@ void init_filter( const Config& conf, { if( sigma0 > 2.0 ) { - cerr << __FILE__ << ":" << __LINE__ << ", ERROR: " - << " Sigma > 2.0 is not supported. Re-size __constant__ array and recompile." - << endl; - exit( -__LINE__ ); + stringstream ss; + ss << "ERROR: " + << " Sigma > 2.0 is not supported. Re-size __constant__ array and recompile."; + POP_FATAL(ss.str()); } if( levels > GAUSS_LEVELS ) { - cerr << __FILE__ << ":" << __LINE__ << ", ERROR: " - << " More than " << GAUSS_LEVELS << " levels not supported. Re-size __constant__ array and recompile." - << endl; - exit( -__LINE__ ); + stringstream ss; + ss << "ERROR: " + << " More than " << GAUSS_LEVELS << " levels not supported. Re-size __constant__ array and recompile."; + POP_FATAL(ss.str()); } if( conf.ifPrintGaussTables() ) { @@ -291,10 +291,9 @@ int GaussInfo::getSpan( float sigma ) const case Config::Fixed15 : return 8; default : - cerr << __FILE__ << ":" << __LINE__ << ", ERROR: " - << " The mode for computing Gauss filter scan is invalid" - << endl; - exit( -__LINE__ ); + stringstream ss; + ss << "ERROR: The mode for computing Gauss filter scan is invalid"; + POP_FATAL(ss.str()); } } diff --git a/src/popsift/popsift.cpp b/src/popsift/popsift.cpp index 253af961..f7e983aa 100755 --- a/src/popsift/popsift.cpp +++ b/src/popsift/popsift.cpp @@ -12,10 +12,13 @@ #include "gauss_filter.h" #include "sift_config.h" #include "sift_pyramid.h" +#include "common/debug_macros.h" #include #include #include +#include +#include using namespace std; @@ -154,7 +157,7 @@ void PopSift::uninit( ) { if(!_isInit) { - std::cout << "[warning] Attempt to release resources from an uninitialized instance" << std::endl; + std::cerr << "[warning] Attempt to release resources from an uninitialized instance" << std::endl; return; } _pipe.uninit(); @@ -248,9 +251,10 @@ SiftJob* PopSift::enqueue( int w, { if( _image_mode != ByteImages ) { - cerr << __FILE__ << ":" << __LINE__ << " Image mode error" << endl - << "E Cannot load byte images into a PopSift pipeline configured for float images" << endl; - exit( -1 ); + stringstream ss; + ss << "Image mode error" << endl + << "E Cannot load byte images into a PopSift pipeline configured for float images"; + POP_FATAL(ss.str()); } AllocTest a = testTextureFit( w, h ); @@ -272,9 +276,10 @@ SiftJob* PopSift::enqueue( int w, { if( _image_mode != FloatImages ) { - cerr << __FILE__ << ":" << __LINE__ << " Image mode error" << endl - << "E Cannot load float images into a PopSift pipeline configured for byte images" << endl; - exit( -1 ); + stringstream ss; + ss << "Image mode error" << endl + << "E Cannot load float images into a PopSift pipeline configured for byte images"; + POP_FATAL(ss.str()); } AllocTest a = testTextureFit( w, h ); @@ -352,20 +357,29 @@ void PopSift::matchPrepareLoop( ) SiftJob* job; while( ( job = p._queue_stage2.pull() ) != nullptr ) { - applyConfiguration(); - - popsift::ImageBase* img = job->getImg(); + popsift::FeaturesDev* features; + try + { + applyConfiguration(); - private_init( img->getWidth(), img->getHeight() ); + popsift::ImageBase* img = job->getImg(); - p._pyramid->step1( _config, img ); - p._unused.push( img ); // uploaded input image no longer needed, release for reuse + private_init(img->getWidth(), img->getHeight()); - p._pyramid->step2( _config ); + p._pyramid->step1(_config, img); + p._unused.push(img); // uploaded input image no longer needed, release for reuse - popsift::FeaturesDev* features = p._pyramid->clone_device_descriptors( _config ); + p._pyramid->step2(_config); - cudaDeviceSynchronize(); + features = p._pyramid->clone_device_descriptors(_config); + cudaDeviceSynchronize(); + } + catch(const std::exception& e) + { + job->setError(std::current_exception()); + job->setFeatures(nullptr); + break; + } job->setFeatures( features ); } @@ -387,9 +401,10 @@ SiftJob::SiftJob( int w, int h, const unsigned char* imageData ) } else { - cerr << __FILE__ << ":" << __LINE__ << " Memory limitation" << endl - << "E Failed to allocate memory for SiftJob" << endl; - exit( -1 ); + stringstream ss; + ss << "Memory limitation" << endl + << "E Failed to allocate memory for SiftJob"; + POP_FATAL(ss.str()); } } @@ -407,9 +422,10 @@ SiftJob::SiftJob( int w, int h, const float* imageData ) } else { - cerr << __FILE__ << ":" << __LINE__ << " Memory limitation" << endl - << "E Failed to allocate memory for SiftJob" << endl; - exit( -1 ); + stringstream ss; + ss << "Memory limitation" << endl + << "E Failed to allocate memory for SiftJob"; + POP_FATAL(ss.str()); } } @@ -458,7 +474,16 @@ popsift::FeaturesHost* SiftJob::getHost() popsift::FeaturesDev* SiftJob::getDev() { - return dynamic_cast( _f.get() ); + popsift::FeaturesBase* features = _f.get(); + if(this->_err != nullptr) { + std::rethrow_exception(this->_err); + } + return dynamic_cast(features); +} + +void SiftJob::setError(std::exception_ptr ptr) +{ + this->_err = ptr; } void PopSift::Pipe::uninit() diff --git a/src/popsift/popsift.h b/src/popsift/popsift.h index e8e83872..3b5f72b8 100755 --- a/src/popsift/popsift.h +++ b/src/popsift/popsift.h @@ -15,9 +15,11 @@ #include +#include #include #include #include +#include #include #include @@ -47,6 +49,7 @@ class SiftJob int _h; unsigned char* _imageData; popsift::ImageBase* _img; + std::exception_ptr _err; #if POPSIFT_IS_DEFINED(POPSIFT_USE_NVTX) nvtxRangeId_t _nvtx_id; #endif @@ -92,6 +95,8 @@ class SiftJob /** fulfill the promise */ void setFeatures( popsift::FeaturesBase* f ); + + void setError(std::exception_ptr ptr); }; /** diff --git a/src/popsift/s_desc_notile.cu b/src/popsift/s_desc_notile.cu index a336898b..9ba8a927 100644 --- a/src/popsift/s_desc_notile.cu +++ b/src/popsift/s_desc_notile.cu @@ -13,6 +13,7 @@ #include #include +#include // 1 -> 19.6 on 980 Ti // 2 -> 19.5 on 980 Ti @@ -151,11 +152,7 @@ bool start_ext_desc_notile( int octave, Octave& oct_obj ) oct_obj.getDataTexLinear( ).tex ); cudaDeviceSynchronize(); cudaError_t err = cudaGetLastError( ); - if( err != cudaSuccess ) { - std::cerr << __FILE__ << ":" << __LINE__ << std::endl - << " cudaGetLastError failed: " << cudaGetErrorString(err) << std::endl; - exit( -__LINE__ ); - } + POP_CUDA_FATAL_TEST(err, "cudaGetLastError failed: "); POP_SYNC_CHK;