From b90cc817931010579b0be223b99632e05aab2d69 Mon Sep 17 00:00:00 2001 From: yboucher Date: Tue, 25 Aug 2026 02:11:44 -0700 Subject: [PATCH 1/5] assert footgun fix --- .../mip_heuristics/feasibility_jump/fj_cpu.cu | 8 +++--- cpp/src/utilities/macros.cuh | 25 ++++++++++++------- 2 files changed, 20 insertions(+), 13 deletions(-) diff --git a/cpp/src/mip_heuristics/feasibility_jump/fj_cpu.cu b/cpp/src/mip_heuristics/feasibility_jump/fj_cpu.cu index b789159953..f774431e75 100644 --- a/cpp/src/mip_heuristics/feasibility_jump/fj_cpu.cu +++ b/cpp/src/mip_heuristics/feasibility_jump/fj_cpu.cu @@ -1012,7 +1012,7 @@ static void apply_move(fj_cpu_climber_t& fj_cpu, delta = new_val - old_val; cuopt_assert(isfinite(new_val), "assignment is not finite"); cuopt_assert(isfinite(delta), "applied delta is not finite"); - cuopt_assert((check_variable_within_bounds(fj_cpu, var_idx, new_val)), + cuopt_assert(check_variable_within_bounds(fj_cpu, var_idx, new_val), "assignment not within bounds"); // Update the LHSs of all involved constraints. @@ -1217,7 +1217,7 @@ static thrust::tuple find_mtm_move( } } if (!isfinite(new_val)) continue; - cuopt_assert((check_variable_within_bounds(fj_cpu, var_idx, new_val)), + cuopt_assert(check_variable_within_bounds(fj_cpu, var_idx, new_val), "new_val is not within bounds"); delta = new_val - val; // more permissive tabu in the case of local minima @@ -1264,7 +1264,7 @@ static thrust::tuple find_mtm_move( auto [score, infeasibility] = compute_score(fj_cpu, var_idx, delta); - cuopt_assert((check_variable_within_bounds(fj_cpu, var_idx, new_val)), ""); + cuopt_assert(check_variable_within_bounds(fj_cpu, var_idx, new_val), ""); cuopt_assert(isfinite(delta), ""); if (fj_cpu.view.move_numerically_stable( @@ -1493,7 +1493,7 @@ static void perturb(fj_cpu_climber_t& fj_cpu) val = std::min(std::max(val, lb), ub); } - cuopt_assert((check_variable_within_bounds(fj_cpu, var_idx, val)), + cuopt_assert(check_variable_within_bounds(fj_cpu, var_idx, val), "value is out of bounds"); fj_cpu.h_assignment[var_idx] = val; } diff --git a/cpp/src/utilities/macros.cuh b/cpp/src/utilities/macros.cuh index d36832015a..d232cca028 100644 --- a/cpp/src/utilities/macros.cuh +++ b/cpp/src/utilities/macros.cuh @@ -14,17 +14,24 @@ // 3) heavy #ifdef ASSERT_MODE #include -#define cuopt_assert(val, msg) assert(val&& msg) -#define cuopt_func_call(func) func; -#else -#define cuopt_assert(val, msg) -#define cuopt_func_call(func) ; -#endif +#include -#ifdef BENCHMARK -#define benchmark_call(func) func; +namespace cuopt::detail { +// handle the argument processing through the C++ parser instead of the preprocessor +// since it chokes on colons in template arguments (e.g. cuopt_assert(std::is_same_v, "message")) +// constexpr because otherwise __host__ __device__ is required and it would break pure host builds +template +constexpr bool assert_msg(T&& cond, const char (&)[N]) +{ + return (bool)cond; +} +} // namespace cuopt::detail + +#define cuopt_assert(...) assert(::cuopt::detail::assert_msg(__VA_ARGS__)) +#define cuopt_func_call(func) func; #else -#define benchmark_call(func) ; +#define cuopt_assert(...) +#define cuopt_func_call(func) ; #endif // For CUDA Driver API From d937007227bda0d23dea150fb8bc0498db4f357b Mon Sep 17 00:00:00 2001 From: yboucher Date: Tue, 25 Aug 2026 02:29:38 -0700 Subject: [PATCH 2/5] style --- cpp/src/utilities/macros.cuh | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/cpp/src/utilities/macros.cuh b/cpp/src/utilities/macros.cuh index d232cca028..7b829c1bcf 100644 --- a/cpp/src/utilities/macros.cuh +++ b/cpp/src/utilities/macros.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 */ @@ -18,8 +18,9 @@ namespace cuopt::detail { // handle the argument processing through the C++ parser instead of the preprocessor -// since it chokes on colons in template arguments (e.g. cuopt_assert(std::is_same_v, "message")) -// constexpr because otherwise __host__ __device__ is required and it would break pure host builds +// since it chokes on colons in template arguments (e.g. cuopt_assert(std::is_same_v, +// "message")) constexpr because otherwise __host__ __device__ is required and it would break pure +// host builds template constexpr bool assert_msg(T&& cond, const char (&)[N]) { From 40042aa7d2aaf60d54a55db2265420d464f9eb58 Mon Sep 17 00:00:00 2001 From: yboucher Date: Tue, 25 Aug 2026 02:31:23 -0700 Subject: [PATCH 3/5] style --- cpp/src/utilities/macros.cuh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/src/utilities/macros.cuh b/cpp/src/utilities/macros.cuh index 7b829c1bcf..b2a891ff06 100644 --- a/cpp/src/utilities/macros.cuh +++ b/cpp/src/utilities/macros.cuh @@ -18,9 +18,9 @@ namespace cuopt::detail { // handle the argument processing through the C++ parser instead of the preprocessor -// since it chokes on colons in template arguments (e.g. cuopt_assert(std::is_same_v, -// "message")) constexpr because otherwise __host__ __device__ is required and it would break pure -// host builds +// since it chokes on colons in template arguments. +// (e.g. cuopt_assert(std::is_same_v, "message")). +// constexpr because otherwise __host__ __device__ is required and it would break pure host builds template constexpr bool assert_msg(T&& cond, const char (&)[N]) { From d56cbe2f0fb8bad62550cdc7f1416f0f11ba1a7a Mon Sep 17 00:00:00 2001 From: yboucher Date: Tue, 25 Aug 2026 03:08:40 -0700 Subject: [PATCH 4/5] fix build --- .../feasibility_jump_kernels.cu | 6 +++--- cpp/src/routing/util_kernels/runtime_checks.cu | 17 ++++++++--------- cpp/src/utilities/macros.cuh | 8 +++++++- 3 files changed, 18 insertions(+), 13 deletions(-) diff --git a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cu b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cu index 0efc93b760..441cfcc01f 100644 --- a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cu +++ b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cu @@ -437,7 +437,7 @@ DI bool check_feasibility(const typename fj_t::climber_data_t::view_t& fj_kahan_babushka_neumaier_sum(delta_it + offset_begin, delta_it + offset_end); cuopt_assert(fj.cstr_satisfied(cIdx, lhs), "constraint violated"); } - cuopt_func_call((check_variable_feasibility(fj, check_integer))); + cuopt_func_call(check_variable_feasibility(fj, check_integer)); return true; } @@ -488,7 +488,7 @@ DI bool save_best_solution(typename fj_t::climber_data_t::view_t& fj) if (*fj.best_excess == 0) { *fj.saved_solution_objective = *fj.incumbent_objective; } } - cuopt_func_call((check_variable_feasibility(fj, false))); + cuopt_func_call(check_variable_feasibility(fj, false)); for (i_t i = threadIdx.x; i < fj.pb.n_variables; i += blockDim.x) { fj.best_assignment[i] = fj.incumbent_assignment[i]; } @@ -499,7 +499,7 @@ DI bool save_best_solution(typename fj_t::climber_data_t::view_t& fj) *fj.weighted_violation_score <= *fj.max_cstr_weight * fj.pb.tolerances.absolute_tolerance, "Violated constraint and score mismatch"); bool check_integer = fj.settings->mode != fj_mode_t::ROUNDING; - cuopt_func_call((check_feasibility(fj, check_integer))); + cuopt_func_call(check_feasibility(fj, check_integer)); } // return whether it is an improving local minimum return improving; diff --git a/cpp/src/routing/util_kernels/runtime_checks.cu b/cpp/src/routing/util_kernels/runtime_checks.cu index b03a278a8e..dd81ae0d55 100644 --- a/cpp/src/routing/util_kernels/runtime_checks.cu +++ b/cpp/src/routing/util_kernels/runtime_checks.cu @@ -59,15 +59,14 @@ __global__ void feasibility_check(typename solution_t::view_t node, route.get_node(i + 1), route.vehicle_info(), d_default_weights, 0.001); if (!res) { *solution.sol_found = 0; - cuopt_assert(true, - printf("Failed node:%d f_excess:%f b_excess:%f route:%d, vehicle:%d " - "n_nodes:%d\n", - node.request.info.node(), - node.time_dim.excess_forward, - node.time_dim.excess_backward, - route_id, - vehicle_id, - n_nodes)); + cuopt_func_call(printf("Failed node:%d f_excess:%f b_excess:%f route:%d, vehicle:%d " + "n_nodes:%d\n", + node.request.info.node(), + node.time_dim.excess_forward, + node.time_dim.excess_backward, + route_id, + vehicle_id, + n_nodes)); cuopt_assert(false, "Node should be feasible combine"); return; } diff --git a/cpp/src/utilities/macros.cuh b/cpp/src/utilities/macros.cuh index b2a891ff06..0c43dca7ba 100644 --- a/cpp/src/utilities/macros.cuh +++ b/cpp/src/utilities/macros.cuh @@ -29,12 +29,18 @@ constexpr bool assert_msg(T&& cond, const char (&)[N]) } // namespace cuopt::detail #define cuopt_assert(...) assert(::cuopt::detail::assert_msg(__VA_ARGS__)) -#define cuopt_func_call(func) func; +#define cuopt_func_call(...) __VA_ARGS__; #else #define cuopt_assert(...) #define cuopt_func_call(func) ; #endif +#ifdef BENCHMARK +#define benchmark_call(...) __VA_ARGS__; +#else +#define benchmark_call(func) ; +#endif + // For CUDA Driver API #define CU_CHECK(expr_to_check, err_func) \ do { \ From a6e2c4021b0b4e38f4cbcb69ec5ae434119ec3e4 Mon Sep 17 00:00:00 2001 From: yboucher Date: Tue, 25 Aug 2026 03:40:23 -0700 Subject: [PATCH 5/5] style, fix build --- .../routing/util_kernels/runtime_checks.cu | 19 ++++++++++--------- cpp/src/utilities/macros.cuh | 6 +++--- 2 files changed, 13 insertions(+), 12 deletions(-) diff --git a/cpp/src/routing/util_kernels/runtime_checks.cu b/cpp/src/routing/util_kernels/runtime_checks.cu index dd81ae0d55..b1142f2e04 100644 --- a/cpp/src/routing/util_kernels/runtime_checks.cu +++ b/cpp/src/routing/util_kernels/runtime_checks.cu @@ -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 */ @@ -59,14 +59,15 @@ __global__ void feasibility_check(typename solution_t::view_t node, route.get_node(i + 1), route.vehicle_info(), d_default_weights, 0.001); if (!res) { *solution.sol_found = 0; - cuopt_func_call(printf("Failed node:%d f_excess:%f b_excess:%f route:%d, vehicle:%d " - "n_nodes:%d\n", - node.request.info.node(), - node.time_dim.excess_forward, - node.time_dim.excess_backward, - route_id, - vehicle_id, - n_nodes)); + cuopt_func_call( + printf("Failed node:%d f_excess:%f b_excess:%f route:%d, vehicle:%d " + "n_nodes:%d\n", + node.request.info.node(), + node.time_dim.excess_forward, + node.time_dim.excess_backward, + route_id, + vehicle_id, + n_nodes)); cuopt_assert(false, "Node should be feasible combine"); return; } diff --git a/cpp/src/utilities/macros.cuh b/cpp/src/utilities/macros.cuh index 0c43dca7ba..53eb98ae55 100644 --- a/cpp/src/utilities/macros.cuh +++ b/cpp/src/utilities/macros.cuh @@ -28,17 +28,17 @@ constexpr bool assert_msg(T&& cond, const char (&)[N]) } } // namespace cuopt::detail -#define cuopt_assert(...) assert(::cuopt::detail::assert_msg(__VA_ARGS__)) +#define cuopt_assert(...) assert(::cuopt::detail::assert_msg(__VA_ARGS__)) #define cuopt_func_call(...) __VA_ARGS__; #else #define cuopt_assert(...) -#define cuopt_func_call(func) ; +#define cuopt_func_call(...) ; #endif #ifdef BENCHMARK #define benchmark_call(...) __VA_ARGS__; #else -#define benchmark_call(func) ; +#define benchmark_call(...) ; #endif // For CUDA Driver API