avoid using thrust in header

Author: koide3

include/gtsam_points/cuda/kernels/correspondence.hpp

@@ -0,0 +1,23 @@
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2025  Kenji Koide (k.koide@aist.go.jp)
+
+#pragma once
+
+#ifdef __CUDACC__
+#define GTSAM_POINTS_HOST_DEVICE __host__ __device__
+#else
+#define GTSAM_POINTS_HOST_DEVICE
+#endif
+
+namespace gtsam_points {
+
+/// @brief A pair of source and target point indices representing a correspondence.
+struct Correspondence {
+  GTSAM_POINTS_HOST_DEVICE Correspondence() : source_idx(-1), target_idx(-1) {}
+  GTSAM_POINTS_HOST_DEVICE Correspondence(int source_idx, int target_idx) : source_idx(source_idx), target_idx(target_idx) {}
+
+  int source_idx;
+  int target_idx;
+};
+
+}  // namespace gtsam_points

include/gtsam_points/cuda/kernels/gicp_derivatives.cuh

@@ -6,6 +6,7 @@
 #include <Eigen/Core>
 #include <thrust/device_vector.h>
 
+#include <gtsam_points/cuda/kernels/correspondence.hpp>
 #include <gtsam_points/cuda/kernels/pose.cuh>
 #include <gtsam_points/cuda/kernels/linearized_system.cuh>
 
@@ -24,9 +25,9 @@ struct gicp_derivatives_kernel {
     source_means_ptr(source_means),
     source_covs_ptr(source_covs) {}
 
-  __device__ LinearizedSystem6 operator()(const thrust::pair<int, int>& source_target_correspondence) const {
-    const int source_idx = source_target_correspondence.first;
-    const int target_idx = source_target_correspondence.second;
+  __device__ LinearizedSystem6 operator()(const Correspondence& source_target_correspondence) const {
+    const int source_idx = source_target_correspondence.source_idx;
+    const int target_idx = source_target_correspondence.target_idx;
     if (source_idx < 0 || target_idx < 0) {
       return LinearizedSystem6::zero();
     }
@@ -93,9 +94,9 @@ struct gicp_error_kernel {
     source_means_ptr(source_means),
     source_covs_ptr(source_covs) {}
 
-  __device__ float operator()(const thrust::pair<int, int>& source_target_correspondence) const {
-    const int source_idx = source_target_correspondence.first;
-    const int target_idx = source_target_correspondence.second;
+  __device__ float operator()(const Correspondence& source_target_correspondence) const {
+    const int source_idx = source_target_correspondence.source_idx;
+    const int target_idx = source_target_correspondence.target_idx;
     if (source_idx < 0 || target_idx < 0) {
       return 0.0f;
     }

include/gtsam_points/factors/integrated_gicp_derivatives.cuh

@@ -6,9 +6,9 @@
 #include <memory>
 #include <Eigen/Core>
 #include <Eigen/Geometry>
-#include <thrust/pair.h>
 
 #include <gtsam_points/ann/kdtree_gpu.hpp>
+#include <gtsam_points/cuda/kernels/correspondence.hpp>
 #include <gtsam_points/types/point_cloud_gpu.hpp>
 
 struct CUstream_st;
@@ -78,6 +78,6 @@ private:
 
   int num_inliers;
   int* num_inliers_gpu;
-  thrust::pair<int, int>* source_target_correspondences;
+  Correspondence* source_target_correspondences;
 };
 }  // namespace gtsam_points

src/gtsam_points/factors/integrated_gicp_derivatives_compute.cu

@@ -34,8 +34,8 @@ struct kdtree_correspondence_kernel {
     target_indices(target_indices),
     target_nodes(target_nodes) {}
 
-  __device__ thrust::pair<int, int> operator()(const thrust::pair<int, int>& source_target) const {
-    const int source_idx = source_target.first;
+  __device__ Correspondence operator()(const Correspondence& source_target) const {
+    const int source_idx = source_target.source_idx;
     if (source_idx < 0) {
       return source_target;
     }
@@ -50,7 +50,7 @@ struct kdtree_correspondence_kernel {
 
     const auto [nn_idx, sq_dist] = nn_search(transed_pt);
 
-    return thrust::make_pair(source_idx, static_cast<int>(nn_idx));
+    return Correspondence(source_idx, static_cast<int>(nn_idx));
   }
 
   const Eigen::Isometry3f* linearization_point_ptr;

src/gtsam_points/factors/integrated_gicp_derivatives_inliers.cu

@@ -8,7 +8,6 @@
 #include <thrust/sequence.h>
 #include <thrust/host_vector.h>
 #include <thrust/transform.h>
-#include <thrust/pair.h>
 
 #include <gtsam_points/cuda/check_error.cuh>
 #include <gtsam_points/cuda/kernels/pose.cuh>
@@ -20,7 +19,7 @@ namespace gtsam_points {
 namespace {
 
 struct init_source_target_indices_kernel {
-  __host__ __device__ thrust::pair<int, int> operator()(const int i) const { return thrust::make_pair(i, -1); }
+  __host__ __device__ Correspondence operator()(const int i) const { return Correspondence(i, -1); }
 };
 
 }  // namespace
@@ -35,7 +34,7 @@ void IntegratedGICPDerivatives::reset_inliers(const Eigen::Isometry3f& x, const
     inlier_evaluation_point_gpu = d_x;
 
     check_error << cudaFreeAsync(source_target_correspondences, stream);
-    check_error << cudaMallocAsync(&source_target_correspondences, sizeof(thrust::pair<int, int>) * source->size(), stream);
+    check_error << cudaMallocAsync(&source_target_correspondences, sizeof(Correspondence) * source->size(), stream);
     thrust::transform(
       thrust::cuda::par_nosync.on(stream),
       thrust::make_counting_iterator<int>(0),

src/gtsam_points/factors/integrated_gicp_derivatives_linearize.cu

@@ -39,8 +39,8 @@ struct kdtree_correspondence_kernel {
     target_nodes(target_nodes),
     max_correspondence_distance_sq(max_correspondence_distance_sq) {}
 
-  __device__ thrust::pair<int, int> operator()(const thrust::pair<int, int>& source_target) const {
-    const int source_idx = source_target.first;
+  __device__ Correspondence operator()(const Correspondence& source_target) const {
+    const int source_idx = source_target.source_idx;
     if (source_idx < 0) {
       return source_target;
     }
@@ -56,10 +56,10 @@ struct kdtree_correspondence_kernel {
     const auto [nn_idx, sq_dist] = nn_search(transed_pt);
 
     if (sq_dist > max_correspondence_distance_sq) {
-      return thrust::make_pair(source_idx, -1);
+      return Correspondence(source_idx, -1);
     }
 
-    return thrust::make_pair(source_idx, static_cast<int>(nn_idx));
+    return Correspondence(source_idx, static_cast<int>(nn_idx));
   }
 
   const Eigen::Isometry3f* linearization_point_ptr;

src/gtsam_points/factors/integrated_gicp_factor_gpu.cpp

@@ -120,7 +120,7 @@ size_t IntegratedGICPFactorGPU::memory_usage() const {
 }
 
 size_t IntegratedGICPFactorGPU::memory_usage_gpu() const {
-  return sizeof(Eigen::Isometry3f) + sizeof(int) + sizeof(std::pair<int, int>) * derivatives->get_num_inliers();
+  return sizeof(Eigen::Isometry3f) + sizeof(int) + sizeof(Correspondence) * derivatives->get_num_inliers();
 }
 
 void IntegratedGICPFactorGPU::set_enable_offloading(bool enable) {