Merge branch 'develop' into exception

.github/workflows/continuous-integration.yml

@@ -0,0 +1,69 @@
+name: Continuous Integration
+
+on:
+  push:
+    branches:
+      - master
+      - develop
+    # Skip jobs when only documentation files are changed
+    paths-ignore:
+      - '**.md'
+      - '**.rst'
+      - 'docs/**'
+  pull_request:
+    paths-ignore:
+      - '**.md'
+      - '**.rst'
+      - 'docs/**'
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    strategy:
+      matrix:
+        container: ["alicevision/popsift-deps:cuda10.2-ubuntu18.04", "alicevision/popsift-deps:cuda11.8.0-ubuntu20.04", "alicevision/popsift-deps:cuda12.1.0-ubuntu22.04"]
+        build_tpe: ["Release", "Debug"]
+        exclude:
+        # excludes debug on this one as it has a segmentation fault during the compilation (!)
+        - container: "alicevision/popsift-deps:cuda12.1.0-ubuntu22.04"
+          build_tpe: "Debug"
+
+    container:
+      image: ${{ matrix.container }}
+
+    env:
+      DEPS_INSTALL_DIR: /opt/
+      BUILD_TYPE: ${{ matrix.build_tpe }}
+      CTEST_OUTPUT_ON_FAILURE: 1
+    steps:
+      - uses: actions/checkout@v2
+
+      - name: Prepare File Tree
+        run: |
+          mkdir ./build
+          mkdir ./build_as_3rdparty
+          mkdir ../popsift_install
+
+      - name: Configure CMake
+        working-directory: ./build
+        run: |
+          cmake .. \
+           -DCMAKE_BUILD_TYPE=${BUILD_TYPE} \
+           -DBUILD_SHARED_LIBS:BOOL=ON \
+           -DCMAKE_PREFIX_PATH="${DEPS_INSTALL_DIR}" \
+           -DPopSift_BUILD_DOCS:BOOL=OFF \
+           -DCMAKE_INSTALL_PREFIX:PATH=$PWD/../../popsift_install
+
+      - name: Build
+        working-directory: ./build
+        run: |
+          make -j$(nproc) install
+
+      - name: Build As Third Party
+        working-directory: ./build_as_3rdparty
+        run: |
+          cmake ../src/application \
+           -DBUILD_SHARED_LIBS:BOOL=ON \
+           -DCMAKE_BUILD_TYPE=${BUILD_TYPE} \
+           -DCMAKE_PREFIX_PATH:PATH="$PWD/../../popsift_install;${DEPS_INSTALL_DIR}"
+          make -j$(nproc)

Dockerfile_deps

@@ -7,7 +7,7 @@ LABEL maintainer="AliceVision Team [email protected]"
 # see https://hub.docker.com/r/nvidia/cuda/
 #
 # For example, to create a ubuntu 16.04 with cuda 8.0 for development, use
-# docker build --build-arg CUDA_TAG=8.0 --tag alicevision/popsift-deps:cuda${CUDA_TAG}-ubuntu${OS_TAG} .
+# docker build --build-arg CUDA_TAG=8.0 --tag alicevision/popsift-deps:cuda${CUDA_TAG}-ubuntu${OS_TAG} -f Dockerfile_deps .
 #
 # then execute with nvidia docker (https://github.com/nvidia/nvidia-docker/wiki/Installation-(version-2.0))
 # docker run -it --runtime=nvidia popsift_deps
@@ -32,12 +32,12 @@ RUN apt-get clean && apt-get update && apt-get install -y --no-install-recommend
         libboost-thread-dev \
  && rm -rf /var/lib/apt/lists/*
 
- # Manually install cmake
+# Manually install cmake
 WORKDIR /tmp/cmake
-ENV CMAKE_VERSION=3.17
+ENV CMAKE_VERSION=3.24
 ENV CMAKE_VERSION_FULL=${CMAKE_VERSION}.2
-RUN wget https://cmake.org/files/v3.17/cmake-${CMAKE_VERSION_FULL}.tar.gz && \
-    tar zxvf cmake-${CMAKE_VERSION_FULL}.tar.gz && \
+RUN wget https://cmake.org/files/v${CMAKE_VERSION}/cmake-${CMAKE_VERSION_FULL}.tar.gz && \
+    tar zxf cmake-${CMAKE_VERSION_FULL}.tar.gz && \
     cd cmake-${CMAKE_VERSION_FULL} && \
     ./bootstrap --prefix=/usr/local  -- -DCMAKE_BUILD_TYPE:STRING=Release -DCMAKE_USE_OPENSSL:BOOL=ON && \
     make -j$(nproc) install && \

README.md

@@ -1,7 +1,8 @@
 
 # PopSift
 
-[![CII Best Practices](https://bestpractices.coreinfrastructure.org/projects/3728/badge)](https://bestpractices.coreinfrastructure.org/projects/3728)  [![Codacy Badge](https://api.codacy.com/project/badge/Grade/8b0f7a68bc0d4df2ac89c6e732917caa)](https://app.codacy.com/manual/alicevision/popsift?utm_source=github.com&utm_medium=referral&utm_content=alicevision/popsift&utm_campaign=Badge_Grade_Settings)
+[![CII Best Practices](https://bestpractices.coreinfrastructure.org/projects/3728/badge)](https://bestpractices.coreinfrastructure.org/projects/3728) 
+[![Codacy Badge](https://app.codacy.com/project/badge/Grade/64f9192b53df46b483e7cf5be7e2dddd)](https://app.codacy.com/gh/alicevision/popsift/dashboard?utm_source=gh&utm_medium=referral&utm_content=&utm_campaign=Badge_grade)
 
 PopSift is an open-source implementation of the SIFT algorithm in CUDA.
 PopSift tries to stick as closely as possible to David Lowe's famous paper [1], while extracting features from an image in real-time at least on an NVidia GTX 980 Ti GPU.
@@ -100,10 +101,12 @@ In particular, users can choose to generate results very similar to VLFeat or re
 We acknowledge that there is at least one SIFT implementation that is vastly faster, but it makes considerable sacrifices in terms of accuracy and compatibility.
 
 ## Continuous integration:
-- [![Build Status](https://travis-ci.org/alicevision/popsift.svg?branch=master)](https://travis-ci.org/alicevision/popsift) master branch.
-- [![Build Status](https://travis-ci.org/alicevision/popsift.svg?branch=develop)](https://travis-ci.org/alicevision/popsift) develop branch.
-- [![Build status](https://ci.appveyor.com/api/projects/status/rsm5269hs288c2ji/branch/develop?svg=true)](https://ci.appveyor.com/project/AliceVision/popsift/branch/develop)
-  develop branch.
+
+* ![Continuous Integration](https://github.com/alicevision/popsift/workflows/Continuous%20Integration/badge.svg?branch=master) master branch on Linux.
+
+* ![Continuous Integration](https://github.com/alicevision/popsift/workflows/Continuous%20Integration/badge.svg?branch=develop) develop branch on Linux.
+
+* [![Build status](https://ci.appveyor.com/api/projects/status/rsm5269hs288c2ji/branch/develop?svg=true)](https://ci.appveyor.com/project/AliceVision/popsift/branch/develop) develop branch on Windows.
 
 ## License
 

cmake/ChooseCudaCC.cmake

@@ -65,7 +65,7 @@ function(chooseCudaCC SUPPORTED_CC SUPPORTED_GENCODE_FLAGS)
 
   set(CC_LIST_BY_SYSTEM_PROCESSOR "")
   if(CMAKE_SYSTEM_PROCESSOR IN_LIST OTHER_SUPPORTED_PROCESSORS)
-    list(APPEND CC_LIST_BY_SYSTEM_PROCESSOR "20;21;30;35;50;52;60;61;70;75;80;86")
+    list(APPEND CC_LIST_BY_SYSTEM_PROCESSOR "20;21;30;35;50;52;60;61;70;75;80;86;89;90")
   endif()
   if(CMAKE_SYSTEM_PROCESSOR IN_LIST TEGRA_SUPPORTED_PROCESSORS)
     list(APPEND CC_LIST_BY_SYSTEM_PROCESSOR "32;53;62;72")
@@ -78,10 +78,17 @@ function(chooseCudaCC SUPPORTED_CC SUPPORTED_GENCODE_FLAGS)
   # Default setting of the CUDA CC versions to compile.
   # Shortening the lists saves a lot of compile time.
   #
-  set(CUDA_MIN_CC 20)
-  set(CUDA_MAX_CC 86)
-  if(CUDA_VERSION VERSION_GREATER_EQUAL 11.1)
+
+  # The current version last time this list was updated was CUDA 12.1.
+  if(CUDA_VERSION VERSION_GREATER_EQUAL 12)
+    set(CUDA_MIN_CC 50)
+    set(CUDA_MAX_CC 90)
+  elseif(CUDA_VERSION VERSION_GREATER_EQUAL 11.8)
+    set(CUDA_MIN_CC 35)
+    set(CUDA_MAX_CC 90)
+  elseif(CUDA_VERSION VERSION_GREATER_EQUAL 11.1)
     set(CUDA_MIN_CC 35)
+    set(CUDA_MAX_CC 86)
   elseif(CUDA_VERSION_MAJOR GREATER_EQUAL 11)
     set(CUDA_MIN_CC 35)
     set(CUDA_MAX_CC 80)
@@ -92,8 +99,10 @@ function(chooseCudaCC SUPPORTED_CC SUPPORTED_GENCODE_FLAGS)
     set(CUDA_MIN_CC 30)
     set(CUDA_MAX_CC 72)
   elseif(CUDA_VERSION_MAJOR GREATER_EQUAL 8)
+    set(CUDA_MIN_CC 20)
     set(CUDA_MAX_CC 62)
   elseif(CUDA_VERSION_MAJOR GREATER_EQUAL 7)
+    set(CUDA_MIN_CC 20)
     set(CUDA_MAX_CC 53)
   else()
     message(FATAL_ERROR "We do not support a CUDA SDK below version 7.0")

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

@@ -170,10 +170,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
 };
 
 /*************************************************************
@@ -340,7 +340,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

@@ -187,12 +187,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 );
@@ -219,13 +214,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 " : " ")

src/popsift/s_filtergrid.cu

@@ -19,9 +19,11 @@
 #if ! POPSIFT_IS_DEFINED(POPSIFT_DISABLE_GRID_FILTER)
 
 #include <thrust/copy.h>
+#include <thrust/count.h>
 #include <thrust/device_vector.h>
 #include <thrust/execution_policy.h>
 #include <thrust/host_vector.h>
+#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/sequence.h>
 #include <thrust/sort.h>