diff --git a/src/popsift/common/device_prop.cu b/src/popsift/common/device_prop.cu index 8436a69b..a55821cd 100644 --- a/src/popsift/common/device_prop.cu +++ b/src/popsift/common/device_prop.cu @@ -274,7 +274,7 @@ bool device_prop_t::checkLimit_2DsurfLayered( int& width, int& height, int& laye std::cerr << __FILE__ << ":" << __LINE__ << ": CUDA device " << currentDevice << std::endl << " does not support layered 2D surfaces " << width - << " bytes wide." << endl; + << " pixels wide." << endl; } width = ptr->maxSurface2DLayered[0]; returnSuccess = false; diff --git a/src/popsift/common/device_prop.h b/src/popsift/common/device_prop.h index ed5db2b2..7a0b142d 100644 --- a/src/popsift/common/device_prop.h +++ b/src/popsift/common/device_prop.h @@ -91,8 +91,7 @@ class device_prop_t /** * @brief Check if a request exceeds the current CUDA device's limit in * surface2DLayered dimensions. surface2DLayered is the writable equivalent - * to texture2DLayered, but the width must be given in bytes, not elements. - * Since we use float, images cannot be as wide as expected. + * to texture2DLayered. * @param[in,out] width Desired width of the texture. * @param[in,out] height Desired height of the texture. * @param[in,out] layers Desired depth of the texture. diff --git a/src/popsift/common/plane_2d.h b/src/popsift/common/plane_2d.h index 86d26f5c..e262714f 100644 --- a/src/popsift/common/plane_2d.h +++ b/src/popsift/common/plane_2d.h @@ -168,10 +168,10 @@ template struct PitchPlane2D : public PlaneT PlaneBase::freeHost2D( this->data, mode ); } __host__ __device__ - inline short getPitchInBytes( ) const { return _pitchInBytes; } + inline size_t getPitchInBytes( ) const { return _pitchInBytes; } protected: - int _pitchInBytes; // pitch width in bytes + size_t _pitchInBytes; // pitch width in bytes }; /************************************************************* @@ -338,7 +338,7 @@ template class Plane2D : public PitchPlane2D __host__ __device__ inline short getHeight( ) const { return _rows; } __host__ __device__ - inline short getByteSize( ) const { return this->_pitchInBytes*_rows; } + inline size_t getByteSize( ) const { return this->_pitchInBytes * _rows; } __host__ inline void allocDev( int w, int h ) { _cols = w; diff --git a/src/popsift/popsift.cpp b/src/popsift/popsift.cpp index 253af961..6224eb93 100755 --- a/src/popsift/popsift.cpp +++ b/src/popsift/popsift.cpp @@ -184,12 +184,7 @@ PopSift::AllocTest PopSift::testTextureFit( int width, int height ) */ int depth = _config.levels + 3; - /* Surfaces have a limited width in bytes, not in elements. - * Our DOG pyramid stores 4/byte floats, so me must check for - * that width. - */ - int byteWidth = width * sizeof(float); - retval = _device_properties.checkLimit_2DsurfLayered( byteWidth, + retval = _device_properties.checkLimit_2DsurfLayered( width, height, depth, warn ); @@ -216,13 +211,13 @@ std::string PopSift::testTextureFitErrorString( AllocTest err, int width, int he { const float upscaleFactor = _config.getUpscaleFactor(); const float scaleFactor = 1.0f / powf( 2.0f, -upscaleFactor ); - int w = ceilf( width * scaleFactor ) * sizeof(float); + int w = ceilf( width * scaleFactor ); int h = ceilf( height * scaleFactor ); int d = _config.levels + 3; _device_properties.checkLimit_2DsurfLayered( w, h, d, false ); - w = w / scaleFactor / sizeof(float); + w = w / scaleFactor; h = h / scaleFactor; ostr << "E Cannot use" << (upscaleFactor==1 ? " default " : " ")