Forward-merge release/26.02 into main

cpp/src/dual_simplex/cuts.cpp

@@ -2775,6 +2775,7 @@ void verify_cuts_against_saved_solution(const csr_matrix_t<i_t, f_t>& cuts,
 #ifdef DUAL_SIMPLEX_INSTANTIATE_DOUBLE
 template class cut_pool_t<int, double>;
 template class cut_generation_t<int, double>;
+template class knapsack_generation_t<int, double>;
 template class tableau_equality_t<int, double>;
 template class mixed_integer_rounding_cut_t<int, double>;
 

cpp/src/linear_programming/pdlp_constants.hpp

@@ -14,7 +14,8 @@
 namespace cuopt::linear_programming::detail {
 inline constexpr int block_size = 128;
 
-static std::pair<size_t, size_t> inline kernel_config_from_batch_size(const size_t batch_size)
+[[maybe_unused]] static std::pair<size_t, size_t> inline kernel_config_from_batch_size(
+  const size_t batch_size)
 {
   assert(batch_size > 0 && "Batch size must be greater than 0");
   const size_t block_size = std::min(static_cast<size_t>(256), batch_size);

cpp/src/mip/feasibility_jump/feasibility_jump.cu

@@ -258,13 +258,15 @@ void fj_t<i_t, f_t>::copy_weights(const weight_t<i_t, f_t>& weights,
                     fj_left_weights  = make_span(cstr_left_weights),
                     fj_right_weights = make_span(cstr_right_weights),
                     new_weights      = make_span(weights.cstr_weights)] __device__(i_t idx) {
-                     fj_weights[idx] = idx >= old_size ? 1. : new_weights[idx];
+                     f_t new_weight = idx >= old_size ? 1. : new_weights[idx];
+                     cuopt_assert(isfinite(new_weight), "invalid weight");
+                     cuopt_assert(new_weight >= 0.0, "invalid weight");
+                     new_weight = std::max(new_weight, 0.0);
+
+                     fj_weights[idx] = idx >= old_size ? 1. : new_weight;
                      // TODO: ask Alice how we can manage the previous left,right weights
-                     fj_left_weights[idx]  = idx >= old_size ? 1. : new_weights[idx];
-                     fj_right_weights[idx] = idx >= old_size ? 1. : new_weights[idx];
-                     cuopt_assert(isfinite(fj_weights[idx]), "invalid weight");
-                     cuopt_assert(isfinite(fj_left_weights[idx]), "invalid left weight");
-                     cuopt_assert(isfinite(fj_right_weights[idx]), "invalid right weight");
+                     fj_left_weights[idx]  = idx >= old_size ? 1. : new_weight;
+                     fj_right_weights[idx] = idx >= old_size ? 1. : new_weight;
                    });
   thrust::transform(handle_ptr->get_thrust_policy(),
                     weights.objective_weight.data(),

cpp/src/mip/feasibility_jump/feasibility_jump_impl_common.cuh

@@ -1,6 +1,6 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
@@ -138,14 +138,16 @@ HDI std::pair<f_t, f_t> feas_score_constraint(
     cuopt_assert(isfinite(new_lhs), "");
     cuopt_assert(isfinite(old_slack) && isfinite(new_slack), "");
 
+    cstr_weight = std::max(cstr_weight, 0.0);
+
     f_t cstr_tolerance = fj.get_corrected_tolerance(cstr_idx);
 
     bool old_viol = fj.excess_score(cstr_idx, current_lhs, c_lb, c_ub) < -cstr_tolerance;
     bool new_viol =
       fj.excess_score(cstr_idx, current_lhs + cstr_coeff * delta, c_lb, c_ub) < -cstr_tolerance;
 
-    bool old_sat = old_lhs < rhs + cstr_tolerance;
-    bool new_sat = new_lhs < rhs + cstr_tolerance;
+    bool old_sat = old_lhs <= rhs + cstr_tolerance;
+    bool new_sat = new_lhs <= rhs + cstr_tolerance;
 
     // equality
     if (fj.pb.integer_equal(c_lb, c_ub)) {

cpp/src/mip/feasibility_jump/feasibility_jump_kernels.cu

@@ -1,6 +1,6 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
@@ -119,6 +119,7 @@ DI void update_weights(typename fj_t<i_t, f_t>::climber_data_t::view_t& fj)
     if (threadIdx.x == 0) {
       // DEVICE_LOG_DEBUG("weight of con %d updated to %g, excess %f\n", cstr_idx, new_weight,
       //     curr_excess_score);
+      new_weight                = std::max(new_weight, 0.0);
       fj.cstr_weights[cstr_idx] = max(fj.cstr_weights[cstr_idx], new_weight);
       if (curr_lower_excess < 0.) {
         fj.cstr_left_weights[cstr_idx] = new_weight;

cpp/src/mip/local_search/local_search.cu

@@ -85,6 +85,7 @@ void local_search_t<i_t, f_t>::start_cpufj_scratch_threads(population_t<i_t, f_t
     cpu_fj.fj_cpu->improvement_callback = [&population, problem_ptr = context.problem_ptr](
                                             f_t obj, const std::vector<f_t>& h_vec) {
       population.add_external_solution(h_vec, obj, solution_origin_t::CPUFJ);
+      (void)problem_ptr;
       if (obj < local_search_best_obj) {
         CUOPT_LOG_TRACE("******* New local search best obj %g, best overall %g",
                         problem_ptr->get_user_obj_from_solver_obj(obj),