Merge pull request #35 from alicevision/cuda7filter

workable thrust filter with cuda 7
alicevision · Feb 13, 2018 · b8ddf1b · b8ddf1b
2 parents 8992a05 + 26a544c
commit b8ddf1b
Showing 1 changed file with 6 additions and 16 deletions.
diff --git a/src/popsift/s_filtergrid.cu b/src/popsift/s_filtergrid.cu
@@ -17,7 +17,7 @@
 
 using namespace std;
 
-#if (__CUDACC_VER__ >= 80000) && not defined(DISABLE_GRID_FILTER)
+#if not defined(DISABLE_GRID_FILTER)
 
 #include <thrust/device_vector.h>
 #include <thrust/sequence.h>
@@ -67,18 +67,6 @@ struct FunctionExtractCell
     }
 };
 
-struct FunctionReversePosition
-{
-    const int _total;
-    FunctionReversePosition( int total ) : _total(total) { }
-
-    __host__ __device__
-    inline int operator()(int val) const
-    {
-        return _total - val - 1;
-    }
-};
-
 struct FunctionIsAbove
 {
     int _limit;
@@ -239,12 +227,14 @@ int Pyramid::extrema_filter_grid( const Config& conf, int ext_total )
     // inclusive prefix sum
     thrust::inclusive_scan( h_cell_counts.begin(), h_cell_counts.end(), cell_count_prefix_sums.begin() );
 
-    FunctionReversePosition fun_reverse_pos( n );
+    thrust::host_vector<int> h_reverse_index(n);
+    thrust::sequence( h_reverse_index.begin(), h_reverse_index.end(),
+                      n-1,
+                      -1 );
 
     // sumup[i] = prefix sum[i] + sum( cell[i] copied into remaining cells )
     thrust::transform( h_cell_counts.begin(), h_cell_counts.end(),
-                       thrust::make_transform_iterator( thrust::make_counting_iterator<int>(0),
-                                                        fun_reverse_pos ),
+                       h_reverse_index.begin(),
                        cell_count_sumup.begin(),
                        thrust::multiplies<int>() );
     thrust::transform( cell_count_sumup.begin(), cell_count_sumup.end(),