Add missing nvtx ranges and cuda checks

cpp/src/dual_simplex/branch_and_bound.cpp

@@ -321,6 +321,7 @@ bool branch_and_bound_t<i_t, f_t>::repair_solution(
   f_t& repaired_obj,
   std::vector<f_t>& repaired_solution) const
 {
+  raft::common::nvtx::range scope("repair_solution");
   bool feasible = false;
   repaired_obj  = std::numeric_limits<f_t>::quiet_NaN();
   i_t n         = original_lp.num_cols;
@@ -399,6 +400,7 @@ branch_and_bound_t<i_t, f_t>::branch_and_bound_t(
 template <typename i_t, typename f_t>
 mip_status_t branch_and_bound_t<i_t, f_t>::solve(mip_solution_t<i_t, f_t>& solution)
 {
+  raft::common::nvtx::range scope("solve_branch_and_bound");
   mip_status_t status = mip_status_t::UNSET;
   mip_solution_t<i_t, f_t> incumbent(original_lp.num_cols);
   if (guess.size() != 0) {

cpp/src/dual_simplex/crossover.cpp

@@ -1032,6 +1032,7 @@ crossover_status_t crossover(const lp_problem_t<i_t, f_t>& lp,
                              lp_solution_t<i_t, f_t>& solution,
                              std::vector<variable_status_t>& vstatus)
 {
+  raft::common::nvtx::range scope("crossover");
   const i_t m         = lp.num_rows;
   const i_t n         = lp.num_cols;
   f_t crossover_start = tic();

cpp/src/dual_simplex/solve.cpp

@@ -115,6 +115,7 @@ lp_status_t solve_linear_program_advanced(const lp_problem_t<i_t, f_t>& original
                                           std::vector<variable_status_t>& vstatus,
                                           std::vector<f_t>& edge_norms)
 {
+  raft::common::nvtx::range scope("solve_linear_program_advanced");
   lp_status_t lp_status = lp_status_t::UNSET;
   lp_problem_t<i_t, f_t> presolved_lp(1, 1, 1);
   presolve_info_t<i_t, f_t> presolve_info;
@@ -241,6 +242,7 @@ lp_status_t solve_linear_program(const user_problem_t<i_t, f_t>& user_problem,
                                  const simplex_solver_settings_t<i_t, f_t>& settings,
                                  lp_solution_t<i_t, f_t>& solution)
 {
+  raft::common::nvtx::range scope("solve_linear_program");
   f_t start_time = tic();
   lp_problem_t<i_t, f_t> original_lp(1, 1, 1);
   std::vector<i_t> new_slacks;
@@ -267,6 +269,7 @@ i_t solve(const user_problem_t<i_t, f_t>& problem,
           const simplex_solver_settings_t<i_t, f_t>& settings,
           std::vector<f_t>& primal_solution)
 {
+  raft::common::nvtx::range scope("solve");
   i_t status;
   if (is_mip(problem) && !settings.relaxation) {
     branch_and_bound_t branch_and_bound(problem, settings);
@@ -305,6 +308,7 @@ i_t solve_mip_with_guess(const user_problem_t<i_t, f_t>& problem,
                          const std::vector<f_t>& guess,
                          mip_solution_t<i_t, f_t>& solution)
 {
+  raft::common::nvtx::range scope("solve_mip_with_guess");
   i_t status;
   if (is_mip(problem)) {
     branch_and_bound_t branch_and_bound(problem, settings);

cpp/src/dual_simplex/solve.hpp

@@ -22,6 +22,8 @@
 #include <dual_simplex/simplex_solver_settings.hpp>
 #include <dual_simplex/types.hpp>
 
+#include <raft/core/nvtx.hpp>
+
 namespace cuopt::linear_programming::dual_simplex {
 
 template <typename i_t, typename f_t>

cpp/src/mip/diversity/diversity_manager.cu

@@ -263,6 +263,7 @@ bool diversity_manager_t<i_t, f_t>::run_presolve(f_t time_limit)
 template <typename i_t, typename f_t>
 void diversity_manager_t<i_t, f_t>::generate_quick_feasible_solution()
 {
+  raft::common::nvtx::range fun_scope("generate_quick_feasible_solution");
   solution_t<i_t, f_t> solution(*problem_ptr);
   // min 1 second, max 10 seconds
   const f_t generate_fast_solution_time = min(10., max(1., timer.remaining_time() / 20.));

cpp/src/mip/diversity/population.cu

@@ -320,6 +320,7 @@ i_t population_t<i_t, f_t>::add_solution(solution_t<i_t, f_t>&& sol)
 template <typename i_t, typename f_t>
 void population_t<i_t, f_t>::normalize_weights()
 {
+  raft::common::nvtx::range fun_scope("normalize_weights");
   CUOPT_LOG_DEBUG("Normalizing weights");
 
   rmm::device_scalar<f_t> l2_norm(problem_ptr->handle_ptr->get_stream());

cpp/src/mip/feasibility_jump/feasibility_jump.cu

@@ -397,7 +397,7 @@ void fj_t<i_t, f_t>::climber_init(i_t climber_idx, const rmm::cuda_stream_view&
   climber->best_l1_distance.set_value_to_zero_async(climber_stream);
   climber->weighted_violation_score.set_value_to_zero_async(climber_stream);
   init_lhs_and_violation<i_t, f_t><<<256, 256, 0, climber_stream.value()>>>(view);
-
+  RAFT_CHECK_CUDA(climber_stream);
   // initialize the best_objective values according to the initial assignment
   f_t best_obj = compute_objective_from_vec<i_t, f_t>(
     climber->incumbent_assignment, pb_ptr->objective_coefficients, climber_stream);
@@ -508,6 +508,7 @@ void fj_t<i_t, f_t>::climber_init(i_t climber_idx, const rmm::cuda_stream_view&
     load_balancing_init_cstr_bounds_csr<i_t, f_t><<<4096, 128, 0, climber_stream.value()>>>(
       view, view.row_size_nonbin_prefix_sum, view.work_id_to_nonbin_var_idx);
 
+  RAFT_CHECK_CUDA(climber_stream);
   cuopt_assert(
     pb_ptr->binary_indices.size() + pb_ptr->nonbinary_indices.size() == pb_ptr->n_variables,
     "invalid variable indices total");
@@ -542,6 +543,7 @@ template <typename i_t, typename f_t>
 void fj_t<i_t, f_t>::load_balancing_score_update(const rmm::cuda_stream_view& stream,
                                                  i_t climber_idx)
 {
+  raft::common::nvtx::range scope("load_balancing_score_update");
   auto [grid_load_balancing_prepare, blocks_load_balancing_prepare] =
     load_balancing_prepare_launch_dims;
   auto [grid_load_balancing_binary, blocks_load_balancing_binary] =
@@ -573,6 +575,7 @@ void fj_t<i_t, f_t>::load_balancing_score_update(const rmm::cuda_stream_view& st
                                                      blocks_load_balancing_binary,
                                                      0,
                                                      data.load_balancing_bin_stream.view()>>>(v);
+    RAFT_CHECK_CUDA(data.load_balancing_bin_stream);
     data.load_balancing_bin_finished_event.record(data.load_balancing_bin_stream.view());
   }
   if (pb_ptr->nonbinary_indices.size() > 0) {
@@ -582,10 +585,12 @@ void fj_t<i_t, f_t>::load_balancing_score_update(const rmm::cuda_stream_view& st
          blocks_load_balancing_mtm_compute_candidates,
          0,
          data.load_balancing_nonbin_stream.view()>>>(v);
+    RAFT_CHECK_CUDA(data.load_balancing_nonbin_stream);
     load_balancing_mtm_compute_scores<i_t, f_t><<<grid_load_balancing_mtm_compute_scores,
                                                   blocks_load_balancing_mtm_compute_scores,
                                                   0,
                                                   data.load_balancing_nonbin_stream.view()>>>(v);
+    RAFT_CHECK_CUDA(data.load_balancing_nonbin_stream);
     data.load_balancing_nonbin_finished_event.record(data.load_balancing_nonbin_stream.view());
   }
 
@@ -767,18 +772,21 @@ void fj_t<i_t, f_t>::run_step_device(const rmm::cuda_stream_view& climber_stream
   }
 
   if (use_graph) cudaGraphLaunch(graph_instance, climber_stream);
+  climber_stream.synchronize();
 }
 
 template <typename i_t, typename f_t>
 void fj_t<i_t, f_t>::refresh_lhs_and_violation(const rmm::cuda_stream_view& stream, i_t climber_idx)
 {
+  raft::common::nvtx::range scope("refresh_lhs_and_violation");
   auto& data = *climbers[climber_idx];
   auto v     = data.view();
 
   data.violated_constraints.clear(stream);
   data.violation_score.set_value_to_zero_async(stream);
   data.weighted_violation_score.set_value_to_zero_async(stream);
   init_lhs_and_violation<i_t, f_t><<<4096, 256, 0, stream>>>(v);
+  RAFT_CHECK_CUDA(stream);
 }
 
 template <typename i_t, typename f_t>
@@ -855,6 +863,7 @@ i_t fj_t<i_t, f_t>::host_loop(solution_t<i_t, f_t>& solution, i_t climber_idx)
       i_t iterations = data.iterations.value(climber_stream);
       // make sure we have the current incumbent saved (e.g. in the case of a timeout)
       update_best_solution_kernel<i_t, f_t><<<1, blocks_resetmoves, 0, climber_stream>>>(v);
+      RAFT_CHECK_CUDA(climber_stream);
       // check feasibility with the relative tolerance rather than the violation score
       raft::copy(solution.assignment.data(),
                  data.best_assignment.data(),

cpp/src/mip/local_search/feasibility_pump/feasibility_pump.cu

@@ -146,6 +146,7 @@ bool feasibility_pump_t<i_t, f_t>::linear_project_onto_polytope(solution_t<i_t,
                                                                 f_t ratio_of_set_integers,
                                                                 bool longer_lp_run)
 {
+  raft::common::nvtx::range scope("linear_project_onto_polytope");
   CUOPT_LOG_DEBUG("linear projection of fp");
   auto h_assignment            = solution.get_host_assignment();
   auto h_variable_upper_bounds = cuopt::host_copy(solution.problem_ptr->variable_upper_bounds,
@@ -248,135 +249,11 @@ bool feasibility_pump_t<i_t, f_t>::linear_project_onto_polytope(solution_t<i_t,
   return true;
 }
 
-template <typename i_t, typename f_t>
-bool feasibility_pump_t<i_t, f_t>::random_round_with_fj(solution_t<i_t, f_t>& solution,
-                                                        timer_t& round_timer)
-{
-  const i_t n_tries = 200;
-  bool is_feasible  = false;
-  rmm::device_uvector<f_t> original_assign(solution.assignment, solution.handle_ptr->get_stream());
-  rmm::device_uvector<f_t> best_rounding(solution.assignment, solution.handle_ptr->get_stream());
-  f_t best_fj_qual = std::numeric_limits<f_t>::max();
-  i_t i;
-  for (i = 0; i < n_tries; ++i) {
-    if (round_timer.check_time_limit()) { break; }
-    CUOPT_LOG_DEBUG("Trying random with FJ");
-    is_feasible = solution.round_nearest();
-    if (is_feasible) {
-      CUOPT_LOG_DEBUG("Feasible found after random round");
-      return true;
-    }
-    // run fj_single descent
-    fj.settings.mode = fj_mode_t::GREEDY_DESCENT;
-    // fj.settings.n_of_minimums_for_exit = 1;
-    fj.settings.update_weights  = false;
-    fj.settings.feasibility_run = true;
-    fj.settings.time_limit      = round_timer.remaining_time();
-    i_t original_sync_period    = fj.settings.parameters.sync_period;
-    // iterations_per_graph is 10
-    fj.settings.parameters.sync_period = 500;
-    is_feasible                        = fj.solve(solution);
-    fj.settings.parameters.sync_period = original_sync_period;
-    cuopt_assert(solution.test_number_all_integer(), "All integers must be rounded");
-    if (is_feasible) {
-      CUOPT_LOG_DEBUG("Feasible found after random round FJ run");
-      return true;
-    }
-    f_t current_qual = solution.get_quality(fj.cstr_weights, fj.objective_weight);
-    if (current_qual < best_fj_qual) {
-      CUOPT_LOG_DEBUG("Updating best quality of roundings %f", current_qual);
-      best_fj_qual = current_qual;
-      raft::copy(best_rounding.data(),
-                 solution.assignment.data(),
-                 solution.assignment.size(),
-                 solution.handle_ptr->get_stream());
-    }
-    raft::copy(solution.assignment.data(),
-               original_assign.data(),
-               solution.assignment.size(),
-               solution.handle_ptr->get_stream());
-  }
-  n_fj_single_descents                       = i;
-  const i_t min_sols_to_run_fj               = 1;
-  const i_t non_improving_local_minima_count = 200;
-  // if at least 5 solutions are explored, run longer fj
-  if (i >= min_sols_to_run_fj) {
-    // run longer fj on best sol
-    raft::copy(solution.assignment.data(),
-               best_rounding.data(),
-               solution.assignment.size(),
-               solution.handle_ptr->get_stream());
-    rmm::device_uvector<f_t> original_weights(fj.cstr_weights, solution.handle_ptr->get_stream());
-    fj.settings.mode                   = fj_mode_t::EXIT_NON_IMPROVING;
-    fj.settings.update_weights         = true;
-    fj.settings.feasibility_run        = true;
-    fj.settings.time_limit             = 0.5;
-    fj.settings.n_of_minimums_for_exit = non_improving_local_minima_count;
-    is_feasible                        = fj.solve(solution);
-    // restore the weights
-    raft::copy(fj.cstr_weights.data(),
-               original_weights.data(),
-               fj.cstr_weights.size(),
-               solution.handle_ptr->get_stream());
-  }
-  if (is_feasible) {
-    CUOPT_LOG_DEBUG("Feasible found after %d minima FJ run", non_improving_local_minima_count);
-    return true;
-  }
-  raft::copy(solution.assignment.data(),
-             original_assign.data(),
-             solution.assignment.size(),
-             solution.handle_ptr->get_stream());
-  solution.handle_ptr->sync_stream();
-  return is_feasible;
-}
-
-template <typename i_t, typename f_t>
-bool feasibility_pump_t<i_t, f_t>::round_multiple_points(solution_t<i_t, f_t>& solution)
-{
-  n_fj_single_descents     = 0;
-  const f_t max_time_limit = last_lp_time * 0.1;
-  timer_t round_timer{min(max_time_limit, timer.remaining_time())};
-  bool is_feasible = random_round_with_fj(solution, round_timer);
-  if (is_feasible) {
-    CUOPT_LOG_DEBUG("Feasible found after random round with fj");
-    return true;
-  }
-  timer_t line_segment_timer{min(1., timer.remaining_time())};
-  i_t n_points_to_search  = n_fj_single_descents;
-  bool is_feasibility_run = true;
-  // create a copy, because assignment is changing within kernel and we want a separate point_1
-  rmm::device_uvector<f_t> starting_point(solution.assignment, solution.handle_ptr->get_stream());
-  is_feasible = line_segment_search.search_line_segment(solution,
-                                                        starting_point,
-                                                        lp_optimal_solution,
-                                                        n_points_to_search,
-                                                        is_feasibility_run,
-                                                        line_segment_timer);
-  if (is_feasible) {
-    CUOPT_LOG_DEBUG("Feasible found after line segment");
-    return true;
-  }
-  // lns.config.run_lp_and_loop = false;
-  // timer_t lns_timer(min(last_lp_time * 0.1, timer.remaining_time()));
-  // is_feasible = lns.do_lns(solution, lns_timer);
-  // if (is_feasible) {
-  //   CUOPT_LOG_DEBUG("Feasible found after inevitable feasibility");
-  //   return true;
-  // }
-  // TODO add the solution with the min distance to the population, if population is given
-  is_feasible = solution.round_nearest();
-  if (is_feasible) {
-    CUOPT_LOG_DEBUG("Feasible found after nearest rounding");
-    return true;
-  }
-  return is_feasible;
-}
-
 // round will use inevitable infeasibility while propagating the bounds
 template <typename i_t, typename f_t>
 bool feasibility_pump_t<i_t, f_t>::round(solution_t<i_t, f_t>& solution)
 {
+  raft::common::nvtx::range fun_scope("round");
   bool result;
   CUOPT_LOG_DEBUG("Rounding the point");
   timer_t bounds_prop_timer(min(2., timer.remaining_time()));
@@ -453,6 +330,7 @@ bool feasibility_pump_t<i_t, f_t>::test_fj_feasible(solution_t<i_t, f_t>& soluti
 template <typename i_t, typename f_t>
 bool feasibility_pump_t<i_t, f_t>::handle_cycle(solution_t<i_t, f_t>& solution)
 {
+  raft::common::nvtx::range fun_scope("handle_cycle");
   CUOPT_LOG_DEBUG("running handle cycle");
   bool is_feasible       = false;
   fp_fj_cycle_time_begin = timer.remaining_time();
@@ -560,6 +438,7 @@ void feasibility_pump_t<i_t, f_t>::save_best_excess_solution(solution_t<i_t, f_t
 template <typename i_t, typename f_t>
 void feasibility_pump_t<i_t, f_t>::relax_general_integers(solution_t<i_t, f_t>& solution)
 {
+  raft::common::nvtx::range fun_scope("relax_general_integers");
   orig_variable_types.resize(solution.problem_ptr->n_variables, solution.handle_ptr->get_stream());
 
   auto var_types  = make_span(solution.problem_ptr->variable_types);
@@ -609,6 +488,7 @@ void feasibility_pump_t<i_t, f_t>::revert_relaxation(solution_t<i_t, f_t>& solut
 template <typename i_t, typename f_t>
 bool feasibility_pump_t<i_t, f_t>::run_single_fp_descent(solution_t<i_t, f_t>& solution)
 {
+  raft::common::nvtx::range fun_scope("run_single_fp_descent");
   // start by doing nearest rounding
   solution.round_nearest();
   raft::copy(last_rounding.data(),

cpp/src/mip/local_search/feasibility_pump/feasibility_pump.cuh

@@ -140,8 +140,6 @@ class feasibility_pump_t {
   void reset();
   void resize_vectors(problem_t<i_t, f_t>& problem, const raft::handle_t* handle_ptr);
   void save_best_excess_solution(solution_t<i_t, f_t>& solution);
-  bool random_round_with_fj(solution_t<i_t, f_t>& solution, timer_t& round_timer);
-  bool round_multiple_points(solution_t<i_t, f_t>& solution);
   void relax_general_integers(solution_t<i_t, f_t>& solution);
   void revert_relaxation(solution_t<i_t, f_t>& solution);
   bool test_fj_feasible(solution_t<i_t, f_t>& solution, f_t time_limit);

cpp/src/mip/local_search/line_segment_search/line_segment_search.cu

@@ -55,6 +55,7 @@ bool line_segment_search_t<i_t, f_t>::search_line_segment(solution_t<i_t, f_t>&
                                                           bool is_feasibility_run,
                                                           cuopt::timer_t& timer)
 {
+  raft::common::nvtx::range scope("search_line_segment");
   CUOPT_LOG_DEBUG("Running line segment search");
   cuopt_assert(point_1.size() == point_2.size(), "size mismatch");
   cuopt_assert(point_1.size() == solution.assignment.size(), "size mismatch");
@@ -136,7 +137,9 @@ bool line_segment_search_t<i_t, f_t>::search_line_segment(solution_t<i_t, f_t>&
              best_assignment.data(),
              solution.assignment.size(),
              solution.handle_ptr->get_stream());
-  return solution.compute_feasibility();
+  bool is_feasible = solution.compute_feasibility();
+  solution.handle_ptr->sync_stream();
+  return is_feasible;
 }
 
 #if MIP_INSTANTIATE_FLOAT

cpp/src/mip/local_search/local_search.cu

@@ -38,17 +38,10 @@ local_search_t<i_t, f_t>::local_search_t(mip_solver_context_t<i_t, f_t>& context
     fj_sol_on_lp_opt(context.problem_ptr->n_variables,
                      context.problem_ptr->handle_ptr->get_stream()),
     fj(context),
-    // fj_tree(fj),
     constraint_prop(context),
     lb_constraint_prop(context),
     line_segment_search(fj),
-    fp(context,
-       fj,
-       // fj_tree,
-       constraint_prop,
-       lb_constraint_prop,
-       line_segment_search,
-       lp_optimal_solution_),
+    fp(context, fj, constraint_prop, lb_constraint_prop, line_segment_search, lp_optimal_solution_),
     rng(cuopt::seed_generator::get_seed())
 {
 }
@@ -344,7 +337,7 @@ bool local_search_t<i_t, f_t>::generate_solution(solution_t<i_t, f_t>& solution,
                                                  bool& early_exit,
                                                  f_t time_limit)
 {
-  raft::common::nvtx::range fun_scope("LS FP Loop");
+  raft::common::nvtx::range fun_scope("generate_solution");
 
   timer_t timer(time_limit);
   auto n_vars         = solution.problem_ptr->n_variables;