remove lbvh_knn.inl

src/cupoch/knn/lbvh_knn.cu

@@ -19,17 +19,40 @@
  * IN THE SOFTWARE.
  **/
 #include <Eigen/Core>
+#include <thrust/sequence.h>
+#include <thrust/sort.h>
 #include "cupoch/knn/lbvh_knn.h"
 
 #include <lbvh_index/aabb.cuh>
 #include <lbvh_index/lbvh.cuh>
 #include <lbvh_index/lbvh_kernels.cuh>
+#include <lbvh_index/query_knn_kernels.cuh>
+
 #include "cupoch/utility/eigen.h"
 #include "cupoch/utility/platform.h"
 #include "cupoch/utility/helper.h"
 
 
 namespace {
+
+template<int Dim>
+struct convert_float3_functor {
+    convert_float3_functor() {}
+    __device__
+    float3 operator() (const Eigen::Matrix<float, Dim, 1>& x) {
+        return make_float3(x[0], x[1], x[2]);
+    }
+};
+
+
+__device__ __host__
+lbvh::AABB to_float3_aabb(const cupoch::knn::AABB& aabb) {
+    lbvh::AABB aabb_f3;
+    aabb_f3.min = make_float3(aabb.first[0], aabb.first[1], aabb.first[2]);
+    aabb_f3.max = make_float3(aabb.second[0], aabb.second[1], aabb.second[2]);
+    return aabb_f3;
+}
+
 template<typename T>
 struct convert_float3_and_aabb_functor {
     convert_float3_and_aabb_functor() {}
@@ -142,6 +165,50 @@ bool LinearBoundingVolumeHierarchyKNN::SetRawData(const utility::device_vector<T
     return true;
 }
 
+template <typename InputIterator, int Dim>
+int LinearBoundingVolumeHierarchyKNN::SearchKNN(InputIterator first,
+                                                InputIterator last,
+                                                int knn,
+                                                utility::device_vector<unsigned int> &indices,
+                                                utility::device_vector<float> &distance2) const {
+    size_t num_query = thrust::distance(first, last);
+    utility::device_vector<float3> data_float3(num_query);
+    thrust::transform(first, last, data_float3.begin(), convert_float3_functor<Dim>());
+
+    utility::device_vector<unsigned long long int> morton_codes(num_query);
+    utility::device_vector<unsigned int> sorted_indices(num_query);
+    thrust::sequence(sorted_indices.begin(), sorted_indices.end());
+    if (sort_queries_) {
+        dim3 block_dim, grid_dim;
+        std::tie(block_dim, grid_dim) = utility::SelectBlockGridSizes(num_query);
+        compute_morton_points_kernel<<<grid_dim, block_dim>>>(
+            thrust::raw_pointer_cast(data_float3.data()), to_float3_aabb(extent_), thrust::raw_pointer_cast(morton_codes.data()), num_query);
+        cudaSafeCall(cudaDeviceSynchronize());
+        thrust::sort_by_key(morton_codes.begin(), morton_codes.end(), sorted_indices.begin());
+    }
+
+    dim3 block_dim, grid_dim;
+    std::tie(block_dim, grid_dim) = utility::SelectBlockGridSizes(num_query);
+    indices.resize(num_query * knn, std::numeric_limits<unsigned int>::max());
+    distance2.resize(num_query * knn, std::numeric_limits<float>::max());
+    utility::device_vector<unsigned int> neighbors(num_query, 0);
+
+    query_knn_kernel<<<grid_dim, block_dim>>>(
+        thrust::raw_pointer_cast(nodes_->data()),
+        thrust::raw_pointer_cast(data_float3_.data()),
+        thrust::raw_pointer_cast(sorted_indices_.data()),
+        root_node_index_,
+        std::numeric_limits<float>::max(),
+        thrust::raw_pointer_cast(data_float3.data()),
+        thrust::raw_pointer_cast(sorted_indices.data()),
+        num_query,
+        thrust::raw_pointer_cast(indices.data()),
+        thrust::raw_pointer_cast(distance2.data()),
+        thrust::raw_pointer_cast(neighbors.data()));
+    cudaSafeCall(cudaDeviceSynchronize());
+    return 1;
+}
+
 template bool LinearBoundingVolumeHierarchyKNN::SetRawData<Eigen::Vector3f>(
         const utility::device_vector<Eigen::Vector3f> &data);
 

src/cupoch/knn/lbvh_knn.h

@@ -21,8 +21,6 @@
 #pragma once
 #include <Eigen/Core>
 
-#include "cupoch/knn/lbvh_knn.h"
-
 #include "cupoch/utility/device_vector.h"
 
 namespace lbvh {
@@ -84,4 +82,3 @@ class LinearBoundingVolumeHierarchyKNN {
 }
 }
 
-#include "cupoch/knn/lbvh_knn.inl"