fast_vgicp_cuda_impl.hpp 6.8 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182
  1. #ifndef FAST_GICP_FAST_VGICP_CUDA_IMPL_HPP
  2. #define FAST_GICP_FAST_VGICP_CUDA_IMPL_HPP
  3. #include <atomic>
  4. #include <Eigen/Core>
  5. #include <Eigen/Geometry>
  6. #include <pcl/point_types.h>
  7. #include <pcl/point_cloud.h>
  8. #include <pcl/search/kdtree.h>
  9. #include <pcl/registration/registration.h>
  10. #include <fast_gicp/gicp/fast_vgicp_cuda.hpp>
  11. #include <fast_gicp/cuda/fast_vgicp_cuda.cuh>
  12. #include <thrust/host_vector.h>
  13. #include <thrust/device_vector.h>
  14. namespace fast_gicp {
  15. template<typename PointSource, typename PointTarget>
  16. FastVGICPCuda<PointSource, PointTarget>::FastVGICPCuda() : LsqRegistration<PointSource, PointTarget>() {
  17. this->reg_name_ = "FastVGICPCuda";
  18. k_correspondences_ = 20;
  19. voxel_resolution_ = 1.0;
  20. regularization_method_ = RegularizationMethod::PLANE;
  21. neighbor_search_method_ = NearestNeighborMethod::CPU_PARALLEL_KDTREE;
  22. vgicp_cuda_.reset(new cuda::FastVGICPCudaCore());
  23. vgicp_cuda_->set_resolution(voxel_resolution_);
  24. vgicp_cuda_->set_kernel_params(0.5, 3.0);
  25. }
  26. template<typename PointSource, typename PointTarget>
  27. FastVGICPCuda<PointSource, PointTarget>::~FastVGICPCuda() {}
  28. template<typename PointSource, typename PointTarget>
  29. void FastVGICPCuda<PointSource, PointTarget>::setCorrespondenceRandomness(int k) {}
  30. template<typename PointSource, typename PointTarget>
  31. void FastVGICPCuda<PointSource, PointTarget>::setResolution(double resolution) {
  32. vgicp_cuda_->set_resolution(resolution);
  33. }
  34. template <typename PointSource, typename PointTarget>
  35. void FastVGICPCuda<PointSource, PointTarget>::setKernelWidth(double kernel_width, double max_dist) {
  36. if (max_dist <= 0.0) {
  37. max_dist = kernel_width * 5.0;
  38. }
  39. vgicp_cuda_->set_kernel_params(kernel_width, max_dist);
  40. }
  41. template<typename PointSource, typename PointTarget>
  42. void FastVGICPCuda<PointSource, PointTarget>::setRegularizationMethod(RegularizationMethod method) {
  43. regularization_method_ = method;
  44. }
  45. template <typename PointSource, typename PointTarget>
  46. void FastVGICPCuda<PointSource, PointTarget>::setNeighborSearchMethod(NeighborSearchMethod method, double radius) {
  47. vgicp_cuda_->set_neighbor_search_method(method, radius);
  48. }
  49. template <typename PointSource, typename PointTarget>
  50. void FastVGICPCuda<PointSource, PointTarget>::setNearestNeighborSearchMethod(NearestNeighborMethod method) {
  51. neighbor_search_method_ = method;
  52. }
  53. template<typename PointSource, typename PointTarget>
  54. void FastVGICPCuda<PointSource, PointTarget>::swapSourceAndTarget() {
  55. vgicp_cuda_->swap_source_and_target();
  56. input_.swap(target_);
  57. }
  58. template<typename PointSource, typename PointTarget>
  59. void FastVGICPCuda<PointSource, PointTarget>::clearSource() {
  60. input_.reset();
  61. }
  62. template<typename PointSource, typename PointTarget>
  63. void FastVGICPCuda<PointSource, PointTarget>::clearTarget() {
  64. target_.reset();
  65. }
  66. template<typename PointSource, typename PointTarget>
  67. void FastVGICPCuda<PointSource, PointTarget>::setInputSource(const PointCloudSourceConstPtr& cloud) {
  68. // the input cloud is the same as the previous one
  69. if(cloud == input_) {
  70. return;
  71. }
  72. pcl::Registration<PointSource, PointTarget, Scalar>::setInputSource(cloud);
  73. std::vector<Eigen::Vector3f, Eigen::aligned_allocator<Eigen::Vector3f>> points(cloud->size());
  74. std::transform(cloud->begin(), cloud->end(), points.begin(), [=](const PointSource& pt) { return pt.getVector3fMap(); });
  75. vgicp_cuda_->set_source_cloud(points);
  76. switch(neighbor_search_method_) {
  77. case NearestNeighborMethod::CPU_PARALLEL_KDTREE: {
  78. std::vector<int> neighbors = find_neighbors_parallel_kdtree<PointSource>(k_correspondences_, cloud);
  79. vgicp_cuda_->set_source_neighbors(k_correspondences_, neighbors);
  80. vgicp_cuda_->calculate_source_covariances(regularization_method_);
  81. } break;
  82. case NearestNeighborMethod::GPU_BRUTEFORCE:
  83. vgicp_cuda_->find_source_neighbors(k_correspondences_);
  84. vgicp_cuda_->calculate_source_covariances(regularization_method_);
  85. break;
  86. case NearestNeighborMethod::GPU_RBF_KERNEL:
  87. vgicp_cuda_->calculate_source_covariances_rbf(regularization_method_);
  88. break;
  89. }
  90. }
  91. template<typename PointSource, typename PointTarget>
  92. void FastVGICPCuda<PointSource, PointTarget>::setInputTarget(const PointCloudTargetConstPtr& cloud) {
  93. // the input cloud is the same as the previous one
  94. if(cloud == target_) {
  95. return;
  96. }
  97. pcl::Registration<PointSource, PointTarget, Scalar>::setInputTarget(cloud);
  98. std::vector<Eigen::Vector3f, Eigen::aligned_allocator<Eigen::Vector3f>> points(cloud->size());
  99. std::transform(cloud->begin(), cloud->end(), points.begin(), [=](const PointTarget& pt) { return pt.getVector3fMap(); });
  100. vgicp_cuda_->set_target_cloud(points);
  101. switch(neighbor_search_method_) {
  102. case NearestNeighborMethod::CPU_PARALLEL_KDTREE: {
  103. std::vector<int> neighbors = find_neighbors_parallel_kdtree<PointTarget>(k_correspondences_, cloud);
  104. vgicp_cuda_->set_target_neighbors(k_correspondences_, neighbors);
  105. vgicp_cuda_->calculate_target_covariances(regularization_method_);
  106. } break;
  107. case NearestNeighborMethod::GPU_BRUTEFORCE:
  108. vgicp_cuda_->find_target_neighbors(k_correspondences_);
  109. vgicp_cuda_->calculate_target_covariances(regularization_method_);
  110. break;
  111. case NearestNeighborMethod::GPU_RBF_KERNEL:
  112. vgicp_cuda_->calculate_target_covariances_rbf(regularization_method_);
  113. break;
  114. }
  115. vgicp_cuda_->create_target_voxelmap();
  116. }
  117. template<typename PointSource, typename PointTarget>
  118. void FastVGICPCuda<PointSource, PointTarget>::computeTransformation(PointCloudSource& output, const Matrix4& guess) {
  119. vgicp_cuda_->set_resolution(this->voxel_resolution_);
  120. LsqRegistration<PointSource, PointTarget>::computeTransformation(output, guess);
  121. }
  122. template<typename PointSource, typename PointTarget>
  123. template<typename PointT>
  124. std::vector<int> FastVGICPCuda<PointSource, PointTarget>::find_neighbors_parallel_kdtree(int k, typename pcl::PointCloud<PointT>::ConstPtr cloud) const {
  125. pcl::search::KdTree<PointT> kdtree;
  126. kdtree.setInputCloud(cloud);
  127. std::vector<int> neighbors(cloud->size() * k);
  128. #pragma omp parallel for schedule(guided, 8)
  129. for(int i = 0; i < cloud->size(); i++) {
  130. std::vector<int> k_indices;
  131. std::vector<float> k_sq_distances;
  132. kdtree.nearestKSearch(cloud->at(i), k, k_indices, k_sq_distances);
  133. std::copy(k_indices.begin(), k_indices.end(), neighbors.begin() + i * k);
  134. }
  135. return neighbors;
  136. }
  137. template<typename PointSource, typename PointTarget>
  138. double FastVGICPCuda<PointSource, PointTarget>::linearize(const Eigen::Isometry3d& trans, Eigen::Matrix<double, 6, 6>* H, Eigen::Matrix<double, 6, 1>* b) {
  139. vgicp_cuda_->update_correspondences(trans);
  140. return vgicp_cuda_->compute_error(trans, H, b);
  141. }
  142. template<typename PointSource, typename PointTarget>
  143. double FastVGICPCuda<PointSource, PointTarget>::compute_error(const Eigen::Isometry3d& trans) {
  144. return vgicp_cuda_->compute_error(trans, nullptr, nullptr);
  145. }
  146. } // namespace fast_gicp
  147. #endif