Fix errors with large images

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;

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.

src/popsift/common/plane_2d.h

@@ -168,10 +168,10 @@ template <typename T> struct PitchPlane2D : public PlaneT<T>
         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 <typename T> class Plane2D : public PitchPlane2D<T>
     __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;

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 " : " ")