routing: add distance breaks

cpp/include/cuopt/routing/data_model_view.hpp

@@ -161,6 +161,34 @@ class data_model_view_t {
                          i_t num_break_locations,
                          bool validate_input = true);
 
+  /**
+   * @brief Add a distance-windowed break for a vehicle.
+   *
+   * The solver inserts one break stop per call within the distance
+   * interval [distance_min, distance_max] (measured along the cost matrix).
+   * Call this function multiple times for the same vehicle to model successive
+   * distance cycles (e.g. first stop: [0, 150], second stop: [150, 300]).
+   *
+   * @param vehicle_id         Vehicle to apply the break to.
+   * @param distance_min       Earliest cumulative route distance at which the
+   *                           vehicle may stop.
+   * @param distance_max       Latest cumulative route distance by which the
+   *                           vehicle must have stopped.
+   * @param duration           Service time at the break location (same unit
+   *                           as other service times in the model).
+   * @param break_locations    Device pointer to eligible break location IDs.
+   *                           Pass nullptr to allow any location.
+   * @param num_break_locations Number of entries in break_locations.
+   * @param validate_input     Run input validation. Defaults to true.
+   */
+  void add_distance_break(i_t vehicle_id,
+                          f_t distance_min,
+                          f_t distance_max,
+                          i_t duration,
+                          i_t const* break_locations,
+                          i_t num_break_locations,
+                          bool validate_input = true);
+
   /**
    * @brief During improvement phase the solver only optimizes for the cost.
    * This function is used to select the best solution accross all climbers

cpp/include/cuopt/routing/routing_structures.hpp

@@ -1,6 +1,6 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
@@ -56,18 +56,48 @@ class break_dimension_t {
   i_t const* break_duration_;
 };
 
+/**
+ * @brief A mandatory break a vehicle must take during its route, triggered either by time
+ *        or by cumulative route distance. If @p locations is empty the break may be taken
+ *        anywhere; otherwise it must occur at one of the specified location IDs.
+ */
 template <typename i_t>
 class vehicle_break_t {
  public:
+  /// Time-windowed break: must start within [earliest, latest].
   vehicle_break_t(i_t earliest, i_t latest, i_t duration, raft::device_span<const i_t> locations)
-    : earliest_(earliest), latest_(latest), duration_(duration), locations_(locations)
+    : earliest_(earliest),
+      latest_(latest),
+      duration_(duration),
+      locations_(locations),
+      is_distance_based_(false),
+      distance_min_(0.f),
+      distance_max_(std::numeric_limits<float>::max())
+  {
+  }
+
+  /// Distance-windowed break: must occur within cumulative distance [distance_min, distance_max].
+  vehicle_break_t(float distance_min,
+                  float distance_max,
+                  i_t duration,
+                  raft::device_span<const i_t> locations)
+    : earliest_(0),
+      latest_(std::numeric_limits<i_t>::max()),
+      duration_(duration),
+      locations_(locations),
+      is_distance_based_(true),
+      distance_min_(distance_min),
+      distance_max_(distance_max)
   {
   }
 
   i_t earliest_;
   i_t latest_;
   i_t duration_;
   raft::device_span<const i_t> locations_{};
+  bool is_distance_based_;
+  float distance_min_;
+  float distance_max_;
 };
 
 template <typename i_t, typename f_t>

cpp/src/routing/data_model_view.cu

@@ -14,8 +14,42 @@
 #include <routing/utilities/check_input.hpp>
 #include <unordered_set>
 
+#include <thrust/sort.h>
 #include <thrust/unique.h>
 
+namespace {
+
+/**
+ * @brief Validates that break locations are within the valid range
+ *        of the location matrix and that all entries are unique.
+ */
+template <typename i_t>
+void validate_break_locations(i_t const* locations,
+                              i_t n,
+                              i_t num_locations,
+                              raft::handle_t const* handle)
+{
+  cuopt::cuopt_expects(
+    n >= 0, cuopt::error_type_t::ValidationError, "Number of break locations must be non-negative");
+  if (n == 0) { return; }
+  cuopt::cuopt_expects(locations != nullptr,
+                       cuopt::error_type_t::ValidationError,
+                       "Break locations cannot be null when num_break_locations > 0");
+  cuopt::cuopt_expects(cuopt::routing::detail::check_min_max_values(
+                         locations, n, i_t{0}, num_locations - 1, handle->get_stream()),
+                       cuopt::error_type_t::ValidationError,
+                       "Break locations should be in [0, num_locations) range");
+  rmm::device_uvector<i_t> tmp(n, handle->get_stream());
+  raft::copy(tmp.begin(), locations, n, handle->get_stream());
+  thrust::sort(handle->get_thrust_policy(), tmp.begin(), tmp.end());
+  auto end         = thrust::unique(handle->get_thrust_policy(), tmp.begin(), tmp.end());
+  i_t unique_items = end - tmp.begin();
+  cuopt::cuopt_expects(n == unique_items,
+                       cuopt::error_type_t::ValidationError,
+                       "There should be unique break locations");
+}
+}  // namespace
+
 namespace cuopt {
 namespace routing {
 
@@ -85,7 +119,7 @@ void data_model_view_t<i_t, f_t>::set_break_locations(i_t const* break_locations
       detail::check_min_max_values(
         break_locations, n_break_locations, 0, num_locations_ - 1, handle_ptr_->get_stream()),
       error_type_t::ValidationError,
-      "Break locations should be at the end of the matrix");
+      "Break locations must be within [0, num_locations)");
     rmm::device_uvector<i_t> tmp_break_nodes(n_break_locations, handle_ptr_->get_stream());
     raft::copy(
       tmp_break_nodes.begin(), break_locations, n_break_locations, handle_ptr_->get_stream());
@@ -155,28 +189,54 @@ void data_model_view_t<i_t, f_t>::add_vehicle_break(i_t vehicle_id,
                                                     i_t num_break_locations,
                                                     bool validate_input)
 {
+  cuopt_expects(0 <= vehicle_id && vehicle_id < fleet_size_,
+                error_type_t::ValidationError,
+                "vehicle_id must be in [0, fleet_size)");
+  cuopt_expects(break_earliest <= break_latest,
+                error_type_t::ValidationError,
+                "Break earliest must be less than or equal than break latest!");
+  cuopt_expects(
+    break_duration >= 0, error_type_t::ValidationError, "break_duration must be non-negative!");
+
+  if (validate_input) {
+    validate_break_locations(break_locations, num_break_locations, num_locations_, handle_ptr_);
+  }
+
   vehicle_breaks_[vehicle_id].push_back(detail::vehicle_break_t<i_t>(
     break_earliest,
     break_latest,
     break_duration,
     raft::device_span<const i_t>(break_locations, num_break_locations)));
+}
 
-  if (validate_input && num_break_locations > 0) {
-    cuopt_expects(
-      detail::check_min_max_values(
-        break_locations, num_break_locations, 0, num_locations_ - 1, handle_ptr_->get_stream()),
-      error_type_t::ValidationError,
-      "Break locations should be at the end of the matrix");
-    rmm::device_uvector<i_t> tmp_break_nodes(num_break_locations, handle_ptr_->get_stream());
-    raft::copy(
-      tmp_break_nodes.begin(), break_locations, num_break_locations, handle_ptr_->get_stream());
-    auto end = thrust::unique(
-      handle_ptr_->get_thrust_policy(), tmp_break_nodes.begin(), tmp_break_nodes.end());
-    i_t unique_items = end - tmp_break_nodes.begin();
-    cuopt_expects(num_break_locations == unique_items,
-                  error_type_t::ValidationError,
-                  "There should be unique break locations");
+template <typename i_t, typename f_t>
+void data_model_view_t<i_t, f_t>::add_distance_break(i_t vehicle_id,
+                                                     f_t distance_min,
+                                                     f_t distance_max,
+                                                     i_t duration,
+                                                     i_t const* break_locations,
+                                                     i_t num_break_locations,
+                                                     bool validate_input)
+{
+  cuopt_expects(0 <= vehicle_id && vehicle_id < fleet_size_,
+                error_type_t::ValidationError,
+                "vehicle_id must be in [0, fleet_size)");
+  cuopt_expects(
+    distance_min >= 0, error_type_t::ValidationError, "distance_min must be non-negative!");
+  cuopt_expects(distance_max > distance_min,
+                error_type_t::ValidationError,
+                "distance break distance_max must be greater than distance_min!");
+  cuopt_expects(duration >= 0, error_type_t::ValidationError, "duration must be non-negative!");
+
+  if (validate_input) {
+    validate_break_locations(break_locations, num_break_locations, num_locations_, handle_ptr_);
   }
+
+  vehicle_breaks_[vehicle_id].push_back(detail::vehicle_break_t<i_t>(
+    static_cast<float>(distance_min),
+    static_cast<float>(distance_max),
+    duration,
+    raft::device_span<const i_t>(break_locations, num_break_locations)));
 }
 
 template <typename i_t, typename f_t>

cpp/src/routing/dimensions.cuh

@@ -1,6 +1,6 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
@@ -182,8 +182,9 @@ using infeasible_cost_t = static_vec_t<dim_t>;
 using objective_cost_t  = static_vec_t<objective_t>;
 
 struct cost_dimension_info_t {
-  bool has_max_constraint = false;
-  HDI bool has_constraints() const { return has_max_constraint; }
+  bool has_max_constraint  = false;
+  bool has_distance_window = false;
+  HDI bool has_constraints() const { return has_max_constraint || has_distance_window; }
 };
 
 struct time_dimension_info_t {
@@ -367,6 +368,14 @@ class enabled_dimensions_t {
    */
   HDI bool has_dimension(dim_t dim) const { return hash & (1 << (int)dim); }
 
+  /// True if any dimension contributing forward/backward window excess is enabled
+  /// (TIME, or DIST when a distance window is configured).
+  HDI bool has_window_dimension() const
+  {
+    return has_dimension(dim_t::TIME) ||
+           (has_dimension(dim_t::DIST) && distance_dim.has_distance_window);
+  }
+
   HDI bool has_objective(objective_t obj) const { return obj_hash & (1 << (int)obj); }
 
   template <size_t I>

cpp/src/routing/ges/compute_delivery_insertions.cuh

@@ -1,6 +1,6 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
@@ -121,15 +121,22 @@ DI bool find_node_insertion(const typename route_t<i_t, f_t, REQUEST>::view_t& c
     if constexpr (insert_mode == insert_mode_t::GES) {
       if (node.forward_feasible(curr_route.vehicle_info()) &&
           node_t<i_t, f_t, REQUEST>::feasible_time_combine(
+            node, curr_route.get_node(node_insertion_idx + 1), curr_route.vehicle_info()) &&
+          node_t<i_t, f_t, REQUEST>::feasible_dist_combine(
             node, curr_route.get_node(node_insertion_idx + 1), curr_route.vehicle_info())) {
         return true;
       }
     }
 
     if constexpr (insert_mode == insert_mode_t::LOCAL_SEARCH) {
-      if (node.time_dim.forward_feasible(
-            curr_route.vehicle_info(), weights[dim_t::TIME], excess_limit) &&
+      if (node_t<i_t, f_t, REQUEST>::window_forward_feasible(
+            node, curr_route.vehicle_info(), weights, excess_limit) &&
           node_t<i_t, f_t, REQUEST>::time_combine(node,
+                                                  curr_route.get_node(node_insertion_idx + 1),
+                                                  curr_route.vehicle_info(),
+                                                  weights,
+                                                  excess_limit) &&
+          node_t<i_t, f_t, REQUEST>::dist_combine(node,
                                                   curr_route.get_node(node_insertion_idx + 1),
                                                   curr_route.vehicle_info(),
                                                   weights,

cpp/src/routing/ges/lexicographic_search/lexicographic_search.cu

@@ -310,10 +310,8 @@ __global__ void lexicographic_search(typename solution_t<i_t, f_t, REQUEST>::vie
           is_forward_feasible = node_stack.advance_insertion(temp_node);
           if (is_forward_feasible) {
             bool is_delivery_time_dist_feasible =
-              node_stack.delivery_node.time_dim.forward_feasible(
-                node_stack.s_route.vehicle_info()) &&
-              node_stack.delivery_node.distance_dim.forward_feasible(
-                node_stack.s_route.vehicle_info());
+              node_t<i_t, f_t, REQUEST>::window_forward_feasible(node_stack.delivery_node,
+                                                                 node_stack.s_route.vehicle_info());
             if (!is_delivery_time_dist_feasible) {
               if (--node_stack.stack_top <= 1) { break; }
               cuopt_assert(node_stack.template k_max_ejection_check<REQUEST>(), "");
@@ -391,10 +389,8 @@ __global__ void lexicographic_search(typename solution_t<i_t, f_t, REQUEST>::vie
           node_stack.top() = temp_node;
           if (forward_feasible) {
             bool is_delivery_time_dist_feasible =
-              node_stack.delivery_node.time_dim.forward_feasible(
-                node_stack.s_route.vehicle_info()) &&
-              node_stack.delivery_node.distance_dim.forward_feasible(
-                node_stack.s_route.vehicle_info());
+              node_t<i_t, f_t, REQUEST>::window_forward_feasible(node_stack.delivery_node,
+                                                                 node_stack.s_route.vehicle_info());
             if (!is_delivery_time_dist_feasible) {
               if (--node_stack.stack_top <= 1) { break; }
               advance = true;

cpp/src/routing/ges/squeeze.cuh

@@ -1,6 +1,6 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
@@ -289,7 +289,13 @@ __global__ void execute_move(typename solution_t<i_t, f_t, REQUEST>::view_t solu
                  abs(orginal_route.template get_dim<dim_t::TIME>()
                        .excess_forward[orginal_route.get_num_nodes()] -
                      orginal_route.template get_dim<dim_t::TIME>().excess_backward[0]) < 0.01,
-               "Backward forward mismatch!");
+               "Time backward/forward mismatch!");
+  cuopt_assert(!(orginal_route.dimensions_info().has_dimension(dim_t::DIST) &&
+                 orginal_route.dimensions_info().distance_dim.has_distance_window) ||
+                 abs(orginal_route.template get_dim<dim_t::DIST>()
+                       .excess_forward[orginal_route.get_num_nodes()] -
+                     orginal_route.template get_dim<dim_t::DIST>().excess_backward[0]) < 0.01,
+               "Distance backward/forward mismatch!");
 }
 
 template <typename i_t, typename f_t, request_t REQUEST, bool squeeze_mode>
@@ -402,6 +408,10 @@ __global__ void eject_inserted_requests(
   }
 }
 
+/**
+ * @brief Inserts each missing break dimension into routes that lack it, one break per outer
+ *        iteration, picking the least-cost insertion position for each. One block per route.
+ */
 template <typename i_t, typename f_t, request_t REQUEST>
 __global__ void squeeze_breaks_kernel(typename solution_t<i_t, f_t, REQUEST>::view_t solution,
                                       const bool include_objective,
@@ -479,7 +489,8 @@ __global__ void squeeze_breaks_kernel(typename solution_t<i_t, f_t, REQUEST>::vi
     __shared__ double reduction_buf[2 * raft::WarpSize];
     block_reduce_ranked(thread_best_cost, t_id, reduction_buf, &reduction_idx);
 
-    if (threadIdx.x == reduction_idx) {
+    if (threadIdx.x == reduction_idx && thread_best_break_node_id >= 0 &&
+        reduction_buf[0] != std::numeric_limits<double>::max()) {
       auto break_node = create_break_node<i_t, f_t, REQUEST>(
         break_nodes, thread_best_break_node_id, solution.problem.dimensions_info);
       // do not update the intra indices yet

cpp/src/routing/local_search/breaks_insertion.cu

@@ -1,6 +1,6 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
@@ -15,6 +15,11 @@ namespace cuopt {
 namespace routing {
 namespace detail {
 
+/**
+ * @brief Looks for a cost-reducing relocation of an existing break node within its route by
+ *        evaluating every alternative position and break-location choice for the same break
+ *        dimension. One block per (route, break_dimension) pair.
+ */
 template <typename i_t, typename f_t, request_t REQUEST>
 __global__ void find_break_insertions_kernel(
   typename solution_t<i_t, f_t, REQUEST>::view_t solution,