Fix sort calls from harris and orb in CUDA

shehzan10 · shehzan10 · commit 87513e07d169 · 2016-04-21T13:42:54.000-04:00
diff --git a/src/backend/cuda/kernel/harris.hpp b/src/backend/cuda/kernel/harris.hpp
@@ -19,6 +19,7 @@
 #include "convolve.hpp"
 #include "gradient.hpp"
 #include "sort_by_key.hpp"
+#include "range.hpp"
 
 namespace cuda
 {
@@ -336,7 +337,9 @@ void harris(unsigned* corners_out,
 
         int sort_elem = harris_responses.strides[3] * harris_responses.dims[3];
         harris_responses.ptr = d_resp_corners;
+        // Create indices using range
         harris_idx.ptr = memAlloc<unsigned>(sort_elem);
+        kernel::range<uint>(harris_idx, 0);
 
         // Sort Harris responses
         sort0ByKey<float, uint, false>(harris_responses, harris_idx);
diff --git a/src/backend/cuda/kernel/orb.hpp b/src/backend/cuda/kernel/orb.hpp
@@ -17,6 +17,7 @@
 #include "convolve.hpp"
 #include "orb_patch.hpp"
 #include "sort_by_key.hpp"
+#include "range.hpp"
 
 #include <boost/scoped_ptr.hpp>
 
@@ -394,10 +395,12 @@ void orb(unsigned* out_feat,
 
         int sort_elem = harris_sorted.strides[3] * harris_sorted.dims[3];
         harris_sorted.ptr = d_score_harris;
+        // Create indices using range
         harris_idx.ptr = memAlloc<unsigned>(sort_elem);
+        kernel::range<uint>(harris_idx, 0);
 
         // Sort features according to Harris responses
-        sort0ByKey<float, uint, false>(harris_sorted, harris_idx);
+        kernel::sort0ByKey<float, uint, false>(harris_sorted, harris_idx);
 
         feat_pyr[i] = std::min(feat_pyr[i], lvl_best[i]);
 
diff --git a/src/backend/cuda/kernel/range.hpp b/src/backend/cuda/kernel/range.hpp
@@ -18,10 +18,10 @@ namespace cuda
     namespace kernel
     {
         // Kernel Launch Config Values
-        static const unsigned TX = 32;
-        static const unsigned TY = 8;
-        static const unsigned TILEX = 512;
-        static const unsigned TILEY = 32;
+        static const unsigned RANGE_TX = 32;
+        static const unsigned RANGE_TY = 8;
+        static const unsigned RANGE_TILEX = 512;
+        static const unsigned RANGE_TILEY = 32;
 
         template<typename T>
         __global__
@@ -74,10 +74,10 @@ namespace cuda
         template<typename T>
         void range(Param<T> out, const int dim)
         {
-            dim3 threads(TX, TY, 1);
+            dim3 threads(RANGE_TX, RANGE_TY, 1);
 
-            int blocksPerMatX = divup(out.dims[0], TILEX);
-            int blocksPerMatY = divup(out.dims[1], TILEY);
+            int blocksPerMatX = divup(out.dims[0], RANGE_TILEX);
+            int blocksPerMatY = divup(out.dims[1], RANGE_TILEY);
             dim3 blocks(blocksPerMatX * out.dims[2],
                         blocksPerMatY * out.dims[3],
                         1);
