From 45bd8ce02d0c2780ae1b8354713f41433b480f04 Mon Sep 17 00:00:00 2001 From: Alice Boucher <160623740+aliceb-nv@users.noreply.github.com> Date: Fri, 6 Feb 2026 17:59:00 +0100 Subject: [PATCH] FJ numerical fix + fix clang build (#835) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit This PR includes a fix for a numerical edge case bug in FJ, along with additional safeguards; changes have been made to ensure compilation works with clang. ## Summary by CodeRabbit * **Bug Fixes** * Improved weight validation and clamping to ensure non-negative values in feasibility algorithms. * Enhanced boundary handling for edge cases in weight copying and updates. * Refined feasibility comparison to use inclusive saturation checks for consistency. * **Chores** * Suppressed compiler warnings and minor formatting adjustments. * Added an explicit template instantiation and updated copyright metadata. Authors: - Alice Boucher (https://github.com/aliceb-nv) Approvers: - Akif ÇÖRDÜK (https://github.com/akifcorduk) URL: https://github.com/NVIDIA/cuopt/pull/835 --- cpp/src/dual_simplex/cuts.cpp | 1 + cpp/src/linear_programming/pdlp_constants.hpp | 3 ++- cpp/src/mip/feasibility_jump/feasibility_jump.cu | 14 ++++++++------ .../feasibility_jump_impl_common.cuh | 8 +++++--- .../feasibility_jump/feasibility_jump_kernels.cu | 3 ++- cpp/src/mip/local_search/local_search.cu | 1 + 6 files changed, 19 insertions(+), 11 deletions(-) diff --git a/cpp/src/dual_simplex/cuts.cpp b/cpp/src/dual_simplex/cuts.cpp index be3f3001d..0bb4e1744 100644 --- a/cpp/src/dual_simplex/cuts.cpp +++ b/cpp/src/dual_simplex/cuts.cpp @@ -2775,6 +2775,7 @@ void verify_cuts_against_saved_solution(const csr_matrix_t& cuts, #ifdef DUAL_SIMPLEX_INSTANTIATE_DOUBLE template class cut_pool_t; template class cut_generation_t; +template class knapsack_generation_t; template class tableau_equality_t; template class mixed_integer_rounding_cut_t; diff --git a/cpp/src/linear_programming/pdlp_constants.hpp b/cpp/src/linear_programming/pdlp_constants.hpp index a2b36a4aa..cf17cc985 100644 --- a/cpp/src/linear_programming/pdlp_constants.hpp +++ b/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 inline kernel_config_from_batch_size(const size_t batch_size) +[[maybe_unused]] static std::pair 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(256), batch_size); diff --git a/cpp/src/mip/feasibility_jump/feasibility_jump.cu b/cpp/src/mip/feasibility_jump/feasibility_jump.cu index c43939cd1..7d721b361 100644 --- a/cpp/src/mip/feasibility_jump/feasibility_jump.cu +++ b/cpp/src/mip/feasibility_jump/feasibility_jump.cu @@ -258,13 +258,15 @@ void fj_t::copy_weights(const weight_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(), diff --git a/cpp/src/mip/feasibility_jump/feasibility_jump_impl_common.cuh b/cpp/src/mip/feasibility_jump/feasibility_jump_impl_common.cuh index fbc5a7b39..e57f0ec9e 100644 --- a/cpp/src/mip/feasibility_jump/feasibility_jump_impl_common.cuh +++ b/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 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)) { diff --git a/cpp/src/mip/feasibility_jump/feasibility_jump_kernels.cu b/cpp/src/mip/feasibility_jump/feasibility_jump_kernels.cu index 5e6fef7e1..41151d377 100644 --- a/cpp/src/mip/feasibility_jump/feasibility_jump_kernels.cu +++ b/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::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; diff --git a/cpp/src/mip/local_search/local_search.cu b/cpp/src/mip/local_search/local_search.cu index 71b944a09..4f56c52ee 100644 --- a/cpp/src/mip/local_search/local_search.cu +++ b/cpp/src/mip/local_search/local_search.cu @@ -85,6 +85,7 @@ void local_search_t::start_cpufj_scratch_threads(population_timprovement_callback = [&population, problem_ptr = context.problem_ptr]( f_t obj, const std::vector& 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),