Fix missing headers for thrust::pair, add trailing return types where needed (#732)

bdice · web-flow · commit 716700bffd1c · 2025-12-18T22:56:47.000Z
This fixes some issues with missing headers and missing return types on device lambdas that are needed for the compiler to correctly infer the return types. This makes cuOpt compatible with CCCL 3.2. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Chris Maes (https://github.com/chris-maes) URL: #732
diff --git a/cpp/src/linear_programming/utilities/problem_checking.cu b/cpp/src/linear_programming/utilities/problem_checking.cu
@@ -45,9 +45,8 @@ void problem_checking_t<i_t, f_t>::check_csr_representation(
   cuopt_expects(thrust::all_of(op_problem.get_handle_ptr()->get_thrust_policy(),
                                op_problem.get_constraint_matrix_indices().cbegin(),
                                op_problem.get_constraint_matrix_indices().cend(),
-                               [n_variables = op_problem.get_n_variables()] __device__(i_t val) {
-                                 return val >= 0 && val < n_variables;
-                               }),
+                               [n_variables = op_problem.get_n_variables()] __device__(
+                                 i_t val) -> bool { return val >= 0 && val < n_variables; }),
                 error_type_t::ValidationError,
                 "A_indices values must positive lower than the number of variables (c size).");
 }
@@ -72,7 +71,7 @@ void problem_checking_t<i_t, f_t>::check_initial_primal_representation(
                                   [lower_bounds = make_span(op_problem.get_variable_lower_bounds()),
                                    upper_bounds = make_span(op_problem.get_variable_upper_bounds()),
                                    assignment_span = make_span(primal_initial_solution),
-                                   int_tol         = 1e-8] __device__(i_t idx) {
+                                   int_tol         = 1e-8] __device__(i_t idx) -> bool {
                                     return assignment_span[idx] < lower_bounds[idx] - int_tol ||
                                            assignment_span[idx] > upper_bounds[idx] + int_tol;
                                   }),
@@ -171,13 +170,14 @@ void problem_checking_t<i_t, f_t>::check_problem_representation(
 
   // Check row type if set
   if (!op_problem.get_row_types().is_empty()) {
-    cuopt_expects(
-      thrust::all_of(op_problem.get_handle_ptr()->get_thrust_policy(),
-                     op_problem.get_row_types().cbegin(),
-                     op_problem.get_row_types().cend(),
-                     [] __device__(char val) { return val == 'E' || val == 'G' || val == 'L'; }),
-      error_type_t::ValidationError,
-      "row_types values must equal to 'E', 'G' or 'L'.");
+    cuopt_expects(thrust::all_of(op_problem.get_handle_ptr()->get_thrust_policy(),
+                                 op_problem.get_row_types().cbegin(),
+                                 op_problem.get_row_types().cend(),
+                                 [] __device__(char val) -> bool {
+                                   return val == 'E' || val == 'G' || val == 'L';
+                                 }),
+                  error_type_t::ValidationError,
+                  "row_types values must equal to 'E', 'G' or 'L'.");
 
     cuopt_expects(
       op_problem.get_row_types().size() == op_problem.get_constraint_bounds().size(),
@@ -322,19 +322,17 @@ bool problem_checking_t<i_t, f_t>::has_crossing_bounds(
     thrust::make_counting_iterator(0),
     thrust::make_counting_iterator(0) + op_problem.get_variable_upper_bounds().size(),
     [upper_bounds = make_span(op_problem.get_variable_upper_bounds()),
-     lower_bounds = make_span(op_problem.get_variable_lower_bounds())] __device__(size_t i) {
-      return upper_bounds[i] >= lower_bounds[i];
-    });
+     lower_bounds = make_span(op_problem.get_variable_lower_bounds())] __device__(size_t i)
+      -> bool { return upper_bounds[i] >= lower_bounds[i]; });
 
   // Check if all constraint bounds are valid (upper >= lower)
   bool all_constraint_bounds_valid = thrust::all_of(
     op_problem.get_handle_ptr()->get_thrust_policy(),
     thrust::make_counting_iterator(0),
     thrust::make_counting_iterator(0) + op_problem.get_constraint_upper_bounds().size(),
     [upper_bounds = make_span(op_problem.get_constraint_upper_bounds()),
-     lower_bounds = make_span(op_problem.get_constraint_lower_bounds())] __device__(size_t i) {
-      return upper_bounds[i] >= lower_bounds[i];
-    });
+     lower_bounds = make_span(op_problem.get_constraint_lower_bounds())] __device__(size_t i)
+      -> bool { return upper_bounds[i] >= lower_bounds[i]; });
 
   // Return true if any bounds are invalid (crossing)
   return !all_variable_bounds_valid || !all_constraint_bounds_valid;
diff --git a/cpp/src/mip/diversity/recombiners/bound_prop_recombiner.cuh b/cpp/src/mip/diversity/recombiners/bound_prop_recombiner.cuh
@@ -9,6 +9,7 @@
 
 #include "recombiner.cuh"
 
+#include <thrust/pair.h>
 #include <mip/local_search/rounding/constraint_prop.cuh>
 #include <mip/relaxed_lp/relaxed_lp.cuh>
 #include <mip/solution/solution.cuh>
diff --git a/cpp/src/mip/feasibility_jump/utils.cuh b/cpp/src/mip/feasibility_jump/utils.cuh
@@ -9,6 +9,7 @@
 
 #include "feasibility_jump.cuh"
 
+#include <thrust/pair.h>
 #include <cuda/atomic>
 #include <raft/core/device_span.hpp>
 #include <rmm/device_scalar.hpp>
diff --git a/cpp/src/mip/local_search/rounding/bounds_repair.cu b/cpp/src/mip/local_search/rounding/bounds_repair.cu
@@ -268,7 +268,7 @@ void bounds_repair_t<i_t, f_t>::compute_damages(problem_t<i_t, f_t>& problem, i_
                       sort_iterator + n_candidates,
                       thrust::make_zip_iterator(thrust::make_tuple(
                         candidates.bound_shift.data(), candidates.variable_index.data())),
-                      [] __device__(auto tuple1, auto tuple2) {
+                      [] __device__(auto tuple1, auto tuple2) -> bool {
                         if (thrust::get<0>(tuple1) < thrust::get<0>(tuple2)) {
                           return true;
                         } else if (thrust::get<0>(tuple1) == thrust::get<0>(tuple2) &&
@@ -291,7 +291,7 @@ i_t bounds_repair_t<i_t, f_t>::find_cutoff_index(const candidates_t<i_t, f_t>& c
     handle_ptr->get_thrust_policy(),
     iterator,
     iterator + n_candidates,
-    [best_cstr_delta, best_damage] __device__(auto tuple) {
+    [best_cstr_delta, best_damage] __device__(auto tuple) -> bool {
       if (thrust::get<0>(tuple) == best_cstr_delta && thrust::get<1>(tuple) <= best_damage) {
         return true;
       }
diff --git a/cpp/src/mip/local_search/rounding/constraint_prop.cuh b/cpp/src/mip/local_search/rounding/constraint_prop.cuh
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <thrust/pair.h>
 #include <mip/local_search/rounding/bounds_repair.cuh>
 #include <mip/presolve/bounds_presolve.cuh>
 #include <mip/presolve/conditional_bound_strengthening.cuh>
diff --git a/cpp/src/mip/presolve/bounds_presolve.cuh b/cpp/src/mip/presolve/bounds_presolve.cuh
@@ -16,6 +16,8 @@
 
 #include <utilities/timer.hpp>
 
+#include <thrust/pair.h>
+
 #include "bounds_update_data.cuh"
 #include "utils.cuh"
 
diff --git a/cpp/src/mip/presolve/bounds_update_helpers.cuh b/cpp/src/mip/presolve/bounds_update_helpers.cuh
@@ -5,6 +5,7 @@
  */
 /* clang-format on */
 
+#include <thrust/pair.h>
 #include <mip/problem/problem.cuh>
 #include <mip/utils.cuh>
 #include "bounds_update_data.cuh"
diff --git a/cpp/src/mip/problem/problem.cu b/cpp/src/mip/problem/problem.cu
@@ -547,20 +547,21 @@ void problem_t<i_t, f_t>::check_problem_representation(bool check_transposed,
   cuopt_expects(thrust::all_of(handle_ptr->get_thrust_policy(),
                                thrust::make_counting_iterator<i_t>(0),
                                thrust::make_counting_iterator<i_t>(n_variables),
-                               [vars_bnd = make_span(variable_bounds)] __device__(i_t idx) {
+                               [vars_bnd = make_span(variable_bounds)] __device__(i_t idx) -> bool {
                                  auto bounds = vars_bnd[idx];
                                  return get_lower(bounds) <= get_upper(bounds);
                                }),
                 error_type_t::ValidationError,
                 "Variable bounds are invalid");
   cuopt_expects(
-    thrust::all_of(handle_ptr->get_thrust_policy(),
-                   thrust::make_counting_iterator<i_t>(0),
-                   thrust::make_counting_iterator<i_t>(n_constraints),
-                   [constraint_lower_bounds = constraint_lower_bounds.data(),
-                    constraint_upper_bounds = constraint_upper_bounds.data()] __device__(i_t idx) {
-                     return constraint_lower_bounds[idx] <= constraint_upper_bounds[idx];
-                   }),
+    thrust::all_of(
+      handle_ptr->get_thrust_policy(),
+      thrust::make_counting_iterator<i_t>(0),
+      thrust::make_counting_iterator<i_t>(n_constraints),
+      [constraint_lower_bounds = constraint_lower_bounds.data(),
+       constraint_upper_bounds = constraint_upper_bounds.data()] __device__(i_t idx) -> bool {
+        return constraint_lower_bounds[idx] <= constraint_upper_bounds[idx];
+      }),
     error_type_t::ValidationError,
     "Constraints bounds are invalid");
 
@@ -582,23 +583,21 @@ void problem_t<i_t, f_t>::check_problem_representation(bool check_transposed,
     cuopt_assert(thrust::all_of(handle_ptr->get_thrust_policy(),
                                 integer_indices.cbegin(),
                                 integer_indices.cend(),
-                                [types = variable_types.data()] __device__(i_t idx) {
+                                [types = variable_types.data()] __device__(i_t idx) -> bool {
                                   return types[idx] == var_t::INTEGER;
                                 }),
                  "The integer indices table contains references to non-integer variables.");
     cuopt_assert(thrust::all_of(handle_ptr->get_thrust_policy(),
                                 binary_indices.cbegin(),
                                 binary_indices.cend(),
-                                [bin_table = is_binary_variable.data()] __device__(i_t idx) {
-                                  return bin_table[idx];
-                                }),
+                                [bin_table = is_binary_variable.data()] __device__(
+                                  i_t idx) -> bool { return bin_table[idx]; }),
                  "The binary indices table contains references to non-binary variables.");
     cuopt_assert(thrust::all_of(handle_ptr->get_thrust_policy(),
                                 nonbinary_indices.cbegin(),
                                 nonbinary_indices.cend(),
-                                [bin_table = is_binary_variable.data()] __device__(i_t idx) {
-                                  return !bin_table[idx];
-                                }),
+                                [bin_table = is_binary_variable.data()] __device__(
+                                  i_t idx) -> bool { return !bin_table[idx]; }),
                  "The non-binary indices table contains references to binary variables.");
     cuopt_assert(
       thrust::all_of(
@@ -607,7 +606,7 @@ void problem_t<i_t, f_t>::check_problem_representation(bool check_transposed,
         thrust::make_counting_iterator<i_t>(n_variables),
         [types     = variable_types.data(),
          bin_table = is_binary_variable.data(),
-         pb_view   = view()] __device__(i_t idx) {
+         pb_view   = view()] __device__(i_t idx) -> bool {
           // ensure the binary variable tables are correct
           if (bin_table[idx]) {
             if (!thrust::binary_search(
@@ -640,7 +639,7 @@ void problem_t<i_t, f_t>::check_problem_representation(bool check_transposed,
         thrust::make_counting_iterator<i_t>(n_variables),
         [types     = variable_types.data(),
          bin_table = is_binary_variable.data(),
-         pb_view   = view()] __device__(i_t idx) {
+         pb_view   = view()] __device__(i_t idx) -> bool {
           // ensure the binary variable tables are correct
           if (bin_table[idx]) {
             if (!thrust::binary_search(
@@ -673,7 +672,7 @@ void problem_t<i_t, f_t>::check_problem_representation(bool check_transposed,
         thrust::make_counting_iterator<i_t>(n_variables),
         [types     = variable_types.data(),
          bin_table = is_binary_variable.data(),
-         pb_view   = view()] __device__(i_t idx) {
+         pb_view   = view()] __device__(i_t idx) -> bool {
           // ensure the binary variable tables are correct
           if (bin_table[idx]) {
             if (!thrust::binary_search(
@@ -1212,7 +1211,7 @@ void problem_t<i_t, f_t>::set_implied_integers(const std::vector<i_t>& implied_i
   objective_is_integral = thrust::all_of(handle_ptr->get_thrust_policy(),
                                          thrust::make_counting_iterator(0),
                                          thrust::make_counting_iterator(n_variables),
-                                         [v = view()] __device__(i_t var_idx) {
+                                         [v = view()] __device__(i_t var_idx) -> bool {
                                            if (v.objective_coefficients[var_idx] == 0) return true;
                                            return v.is_integer(v.objective_coefficients[var_idx]) &&
                                                   (v.variable_types[var_idx] == var_t::INTEGER ||
diff --git a/cpp/src/mip/problem/problem_helpers.cuh b/cpp/src/mip/problem/problem_helpers.cuh
@@ -250,7 +250,7 @@ static void check_csr_representation([[maybe_unused]] const rmm::device_uvector<
   cuopt_assert(thrust::all_of(handle_ptr->get_thrust_policy(),
                               variables.cbegin(),
                               variables.cend(),
-                              [n_variables = n_variables] __device__(i_t val) {
+                              [n_variables = n_variables] __device__(i_t val) -> bool {
                                 return val >= 0 && val < n_variables;
                               }),
                "A_indices values must positive lower than the number of variables (c size).");
@@ -267,7 +267,7 @@ static bool check_var_bounds_sanity(const detail::problem_t<i_t, f_t>& problem)
                    thrust::counting_iterator(0),
                    thrust::counting_iterator((i_t)problem.variable_bounds.size()),
                    [tolerance = problem.tolerances.presolve_absolute_tolerance,
-                    var_bnd   = make_span(problem.variable_bounds)] __device__(i_t index) {
+                    var_bnd   = make_span(problem.variable_bounds)] __device__(i_t index) -> bool {
                      auto var_bounds = var_bnd[index];
                      return (get_lower(var_bounds) > get_upper(var_bounds) + tolerance);
                    });
@@ -283,7 +283,7 @@ static bool check_constraint_bounds_sanity(const detail::problem_t<i_t, f_t>& pr
                    thrust::counting_iterator((i_t)problem.constraint_lower_bounds.size()),
                    [tolerance = problem.tolerances.presolve_absolute_tolerance,
                     lb        = make_span(problem.constraint_lower_bounds),
-                    ub        = make_span(problem.constraint_upper_bounds)] __device__(i_t index) {
+                    ub = make_span(problem.constraint_upper_bounds)] __device__(i_t index) -> bool {
                      return (lb[index] > ub[index] + tolerance);
                    });
   return !crossing_bounds_detected;
diff --git a/cpp/src/mip/solution/solution.cuh b/cpp/src/mip/solution/solution.cuh
@@ -14,6 +14,7 @@
 #include <mip/problem/problem.cuh>
 #include <mip/relaxed_lp/lp_state.cuh>
 
+#include <thrust/pair.h>
 #include <raft/util/cuda_dev_essentials.cuh>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
diff --git a/cpp/src/mip/utils.cuh b/cpp/src/mip/utils.cuh
@@ -188,12 +188,29 @@ inline i_t compute_number_of_integer_var_diff(const rmm::device_uvector<i_t>& in
     thrust::count_if(handle_ptr->get_thrust_policy(),
                      indices.begin(),
                      indices.end(),
-                     [first_ptr, second_ptr, int_tol] __host__ __device__(i_t idx) {
+                     [first_ptr, second_ptr, int_tol] __host__ __device__(i_t idx) -> bool {
                        return integer_equal<f_t>(first_ptr[idx], second_ptr[idx], int_tol);
                      });
   return same_vars;
 }
 
+template <typename i_t, typename f_t>
+struct integer_equal_on_indices_functor {
+  const f_t* first_ptr;
+  const f_t* second_ptr;
+  f_t int_tol;
+
+  integer_equal_on_indices_functor(const f_t* first, const f_t* second, f_t tol)
+    : first_ptr(first), second_ptr(second), int_tol(tol)
+  {
+  }
+
+  __host__ __device__ bool operator()(i_t idx) const
+  {
+    return integer_equal<f_t>(first_ptr[idx], second_ptr[idx], int_tol);
+  }
+};
+
 template <typename i_t, typename f_t>
 bool check_integer_equal_on_indices(const rmm::device_uvector<i_t>& indices,
                                     const rmm::device_uvector<f_t>& first_sol,
@@ -202,14 +219,12 @@ bool check_integer_equal_on_indices(const rmm::device_uvector<i_t>& indices,
                                     const raft::handle_t* handle_ptr)
 {
   cuopt_assert(first_sol.size() == second_sol.size(), "Size mismatch!");
-  auto const first_ptr  = make_span(first_sol);
-  auto const second_ptr = make_span(second_sol);
+  const f_t* first_ptr  = first_sol.data();
+  const f_t* second_ptr = second_sol.data();
   return thrust::all_of(handle_ptr->get_thrust_policy(),
                         indices.begin(),
                         indices.end(),
-                        [first_ptr, second_ptr, int_tol] __host__ __device__(i_t idx) {
-                          return integer_equal<f_t>(first_ptr[idx], second_ptr[idx], int_tol);
-                        });
+                        integer_equal_on_indices_functor<i_t, f_t>(first_ptr, second_ptr, int_tol));
 }
 
 template <typename i_t, typename f_t>
@@ -328,39 +343,64 @@ template <typename f_t>
 bool has_nans(const raft::handle_t* handle_ptr, const rmm::device_uvector<f_t>& vec)
 {
   return thrust::any_of(
-    handle_ptr->get_thrust_policy(), vec.begin(), vec.end(), [] __device__(f_t val) {
+    handle_ptr->get_thrust_policy(), vec.begin(), vec.end(), [] __device__(f_t val) -> bool {
       return isnan(val);
     });
 }
 
+template <typename i_t, typename f_t>
+struct has_integrality_discrepancy_functor {
+  const f_t* assignment_ptr;
+  f_t int_tol;
+
+  has_integrality_discrepancy_functor(const f_t* ptr, f_t tol) : assignment_ptr(ptr), int_tol(tol)
+  {
+  }
+
+  __host__ __device__ bool operator()(i_t idx) const
+  {
+    return !is_integer<f_t>(assignment_ptr[idx], int_tol);
+  }
+};
+
 template <typename i_t, typename f_t>
 bool has_integrality_discrepancy(const raft::handle_t* handle_ptr,
                                  const rmm::device_uvector<i_t>& integer_var_indices,
                                  const rmm::device_uvector<f_t>& assignment,
                                  f_t int_tol)
 {
-  auto const assignment_span = make_span(assignment);
   return thrust::any_of(handle_ptr->get_thrust_policy(),
                         integer_var_indices.begin(),
                         integer_var_indices.end(),
-                        [assignment_span, int_tol] __host__ __device__(i_t idx) {
-                          return !is_integer<f_t>(assignment_span[idx], int_tol);
-                        });
+                        has_integrality_discrepancy_functor<i_t, f_t>(assignment.data(), int_tol));
 }
 
+template <typename i_t, typename f_t>
+struct has_variable_bounds_violation_functor {
+  const f_t* assignment_ptr;
+  typename problem_t<i_t, f_t>::view_t problem_view;
+
+  has_variable_bounds_violation_functor(const f_t* ptr, typename problem_t<i_t, f_t>::view_t view)
+    : assignment_ptr(ptr), problem_view(view)
+  {
+  }
+
+  __device__ bool operator()(i_t idx) const
+  {
+    return !problem_view.check_variable_within_bounds(idx, assignment_ptr[idx]);
+  }
+};
+
 template <typename i_t, typename f_t>
 bool has_variable_bounds_violation(const raft::handle_t* handle_ptr,
                                    const rmm::device_uvector<f_t>& assignment,
                                    problem_t<i_t, f_t>* problem_ptr)
 {
-  auto const assignment_span = make_span(assignment);
-  return thrust::any_of(handle_ptr->get_thrust_policy(),
-                        thrust::make_counting_iterator(0),
-                        thrust::make_counting_iterator(0) + problem_ptr->n_variables,
-                        [assignment_span, problem_view = problem_ptr->view()] __device__(i_t idx) {
-                          return !problem_view.check_variable_within_bounds(idx,
-                                                                            assignment_span[idx]);
-                        });
+  return thrust::any_of(
+    handle_ptr->get_thrust_policy(),
+    thrust::make_counting_iterator(0),
+    thrust::make_counting_iterator(0) + problem_ptr->n_variables,
+    has_variable_bounds_violation_functor<i_t, f_t>(assignment.data(), problem_ptr->view()));
 }
 
 }  // namespace cuopt::linear_programming::detail
diff --git a/cpp/src/routing/local_search/sliding_tsp.cu b/cpp/src/routing/local_search/sliding_tsp.cu
@@ -10,6 +10,7 @@
 #include "../utilities/cuopt_utils.cuh"
 #include "local_search.cuh"
 
+#include <thrust/pair.h>
 #include <cub/cub.cuh>
 
 namespace cuopt {
diff --git a/cpp/src/routing/local_search/vrp/fragment_kernels.cuh b/cpp/src/routing/local_search/vrp/fragment_kernels.cuh
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <thrust/pair.h>
 #include "../../solution/solution.cuh"
 #include "../move_candidates/move_candidates.cuh"
 
diff --git a/cpp/src/routing/solution/route_node_map.cuh b/cpp/src/routing/solution/route_node_map.cuh
diff --git a/cpp/src/routing/utilities/check_input.cu b/cpp/src/routing/utilities/check_input.cu
diff --git a/cpp/src/routing/utilities/md_utils.hpp b/cpp/src/routing/utilities/md_utils.hpp
