diff --git a/cpp/src/mip/diversity/population.cu b/cpp/src/mip/diversity/population.cu index 17c88df3f..0b506cbcb 100644 --- a/cpp/src/mip/diversity/population.cu +++ b/cpp/src/mip/diversity/population.cu @@ -168,10 +168,10 @@ void population_t::add_external_solution(const std::vector& solut external_solution_queue_cpufj.erase(external_solution_queue_cpufj.begin() + worst_obj_idx); } - CUOPT_LOG_INFO("%s added a solution to population, solution queue size %lu with objective %g", - solution_origin_to_string(origin), - external_solution_queue.size(), - problem_ptr->get_user_obj_from_solver_obj(objective)); + CUOPT_LOG_DEBUG("%s added a solution to population, solution queue size %lu with objective %g", + solution_origin_to_string(origin), + external_solution_queue.size(), + problem_ptr->get_user_obj_from_solver_obj(objective)); if (external_solution_queue.size() >= 5) { early_exit_primal_generation = true; } } @@ -227,8 +227,8 @@ std::vector> population_t::get_external_solutions } } if (external_solution_queue.size() > 0) { - CUOPT_LOG_INFO("Consuming B&B solutions, solution queue size %lu", - external_solution_queue.size()); + CUOPT_LOG_DEBUG("Consuming B&B solutions, solution queue size %lu", + external_solution_queue.size()); external_solution_queue.clear(); } external_solution_queue_cpufj.clear(); diff --git a/cpp/src/mip/local_search/rounding/constraint_prop.cu b/cpp/src/mip/local_search/rounding/constraint_prop.cu index 4dfd1b216..76d3916df 100644 --- a/cpp/src/mip/local_search/rounding/constraint_prop.cu +++ b/cpp/src/mip/local_search/rounding/constraint_prop.cu @@ -907,7 +907,7 @@ bool constraint_prop_t::find_integer( CUOPT_LOG_DEBUG("Bounds propagation rounding: unset vars %lu", unset_integer_vars.size()); if (unset_integer_vars.size() == 0) { - CUOPT_LOG_ERROR("No integer variables provided in the bounds prop rounding"); + CUOPT_LOG_DEBUG("No integer variables provided in the bounds prop rounding"); expand_device_copy(orig_sol.assignment, sol.assignment, sol.handle_ptr->get_stream()); cuopt_func_call(orig_sol.test_variable_bounds()); return orig_sol.compute_feasibility(); diff --git a/cpp/src/mip/presolve/spmv_kernels.cuh b/cpp/src/mip/presolve/spmv_kernels.cuh deleted file mode 100644 index 3f11a8f36..000000000 --- a/cpp/src/mip/presolve/spmv_kernels.cuh +++ /dev/null @@ -1,229 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION & AFFILIATES. All rights - * reserved. SPDX-License-Identifier: Apache-2.0 - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include -#include - -namespace cuopt::linear_programming::detail { - -template -__device__ f_t spmv(view_t view, raft::device_span input, i_t tid, i_t beg, i_t end) -{ - f_t out = 0.; - for (i_t i = tid + beg; i < end; i += MAX_EDGE_PER_CNST) { - auto coeff = view.coeff[i]; - auto var = view.elem[i]; - auto in = input[var]; - out += coeff * in; - } - return out; -} - -template -__global__ void lb_spmv_heavy_kernel(i_t id_range_beg, - raft::device_span ids, - raft::device_span pseudo_block_ids, - i_t work_per_block, - view_t view, - raft::device_span input, - raft::device_span tmp_out) -{ - auto idx = ids[blockIdx.x] + id_range_beg; - auto pseudo_block_id = pseudo_block_ids[blockIdx.x]; - i_t item_off_beg = view.offsets[idx] + work_per_block * pseudo_block_id; - i_t item_off_end = min(item_off_beg + work_per_block, view.offsets[idx + 1]); - - typedef cub::BlockReduce BlockReduce; - __shared__ typename BlockReduce::TempStorage temp_storage; - - auto out = spmv(view, input, threadIdx.x, item_off_beg, item_off_end); - - out = BlockReduce(temp_storage).Sum(out); - - if (threadIdx.x == 0) { tmp_out[blockIdx.x] = out; } -} - -template -__global__ void finalize_spmv_kernel(i_t heavy_beg_id, - raft::device_span item_offsets, - raft::device_span tmp_out, - view_t view, - raft::device_span output) -{ - using warp_reduce = cub::WarpReduce; - __shared__ typename warp_reduce::TempStorage temp_storage; - i_t idx = heavy_beg_id + blockIdx.x; - i_t item_idx = view.reorg_ids[idx]; - - i_t item_off_beg = item_offsets[blockIdx.x]; - i_t item_off_end = item_offsets[blockIdx.x + 1]; - f_t out = 0.; - for (i_t i = threadIdx.x + item_off_beg; i < item_off_end; i += blockDim.x) { - out += tmp_out[i]; - } - out = warp_reduce(temp_storage).Sum(out); - if (threadIdx.x == 0) { output[item_idx] = out; } -} - -template -__global__ void lb_spmv_block_kernel(i_t id_range_beg, - view_t view, - raft::device_span input, - raft::device_span output) - -{ - i_t idx = id_range_beg + blockIdx.x; - i_t item_idx = view.reorg_ids[idx]; - i_t item_off_beg = view.offsets[idx]; - i_t item_off_end = view.offsets[idx + 1]; - - typedef cub::BlockReduce BlockReduce; - __shared__ typename BlockReduce::TempStorage temp_storage; - - auto out = spmv(view, input, threadIdx.x, item_off_beg, item_off_end); - - out = BlockReduce(temp_storage).Sum(out); - - if (threadIdx.x == 0) { - // written to old index - output[item_idx] = out; - } -} - -template -__global__ void lb_spmv_sub_warp_kernel(i_t id_range_beg, - i_t id_range_end, - view_t view, - raft::device_span input, - raft::device_span output) -{ - constexpr i_t ids_per_block = BDIM / MAX_EDGE_PER_CNST; - i_t id_beg = blockIdx.x * ids_per_block + id_range_beg; - i_t idx = id_beg + (threadIdx.x / MAX_EDGE_PER_CNST); - i_t item_idx; - if (idx < id_range_end) { item_idx = view.reorg_ids[idx]; } - i_t p_tid = threadIdx.x % MAX_EDGE_PER_CNST; - - i_t head_flag = (p_tid == 0); - - using warp_reduce = cub::WarpReduce; - __shared__ typename warp_reduce::TempStorage temp_storage; - - f_t out = 0.; - - if (idx < id_range_end) { - i_t item_off_beg = view.offsets[idx]; - i_t item_off_end = view.offsets[idx + 1]; - out = spmv(view, input, p_tid, item_off_beg, item_off_end); - } - - out = warp_reduce(temp_storage).Sum(out); - - if (head_flag && (idx < id_range_end)) { output[item_idx] = out; } -} - -#if 1 - -#define BYTE_TO_BINARY(byte) \ - ((byte) & 0x80 ? '1' : '0'), ((byte) & 0x40 ? '1' : '0'), ((byte) & 0x20 ? '1' : '0'), \ - ((byte) & 0x10 ? '1' : '0'), ((byte) & 0x08 ? '1' : '0'), ((byte) & 0x04 ? '1' : '0'), \ - ((byte) & 0x02 ? '1' : '0'), ((byte) & 0x01 ? '1' : '0') - -template -__device__ __forceinline__ void get_sub_warp_bin(i_t* id_warp_beg, - i_t* id_range_end, - i_t* t_p_v, - raft::device_span warp_offsets, - raft::device_span bin_offsets) -{ - i_t warp_id = (blockDim.x * blockIdx.x + threadIdx.x) / 32; - i_t lane_id = threadIdx.x & 31; - bool pred = false; - if (lane_id < warp_offsets.size()) { pred = (warp_id >= warp_offsets[lane_id]); } - unsigned int m = __ballot_sync(0xffffffff, pred); - i_t seg = 31 - __clz(m); - i_t it_per_warp = (1 << (5 - seg)); // item per warp = 32/(2^seg) - if (5 - seg < 0) { - *t_p_v = 0; - return; - } - i_t beg = bin_offsets[seg] + (warp_id - warp_offsets[seg]) * it_per_warp; - i_t end = bin_offsets[seg + 1]; - *id_warp_beg = beg; - *id_range_end = end; - *t_p_v = (1 << seg); -} - -template -__device__ void spmv_sub_warp(i_t id_warp_beg, - i_t id_range_end, - view_t view, - raft::device_span input, - raft::device_span output) -{ - i_t lane_id = (threadIdx.x & 31); - i_t idx = id_warp_beg + (lane_id / MAX_EDGE_PER_CNST); - i_t item_idx; - if (idx < id_range_end) { item_idx = view.reorg_ids[idx]; } - i_t p_tid = lane_id & (MAX_EDGE_PER_CNST - 1); - - i_t head_flag = (p_tid == 0); - - using warp_reduce = cub::WarpReduce; - __shared__ typename warp_reduce::TempStorage temp_storage; - - f_t out = 0.; - - if (idx < id_range_end) { - i_t item_off_beg = view.offsets[idx]; - i_t item_off_end = view.offsets[idx + 1]; - out = spmv(view, input, p_tid, item_off_beg, item_off_end); - } - - out = warp_reduce(temp_storage).Sum(out); - - if (head_flag && (idx < id_range_end)) { output[item_idx] = out; } -} - -template -__global__ void lb_spmv_sub_warp_kernel(view_t view, - raft::device_span input, - raft::device_span output, - raft::device_span warp_item_offsets, - raft::device_span warp_item_id_offsets) -{ - i_t id_warp_beg, id_range_end, t_p_v; - get_sub_warp_bin( - &id_warp_beg, &id_range_end, &t_p_v, warp_item_offsets, warp_item_id_offsets); - - if (t_p_v == 1) { - spmv_sub_warp(id_warp_beg, id_range_end, view, input, output); - } else if (t_p_v == 2) { - spmv_sub_warp(id_warp_beg, id_range_end, view, input, output); - } else if (t_p_v == 4) { - spmv_sub_warp(id_warp_beg, id_range_end, view, input, output); - } else if (t_p_v == 8) { - spmv_sub_warp(id_warp_beg, id_range_end, view, input, output); - } else if (t_p_v == 16) { - spmv_sub_warp(id_warp_beg, id_range_end, view, input, output); - } -} -#endif - -} // namespace cuopt::linear_programming::detail diff --git a/cpp/src/mip/presolve/third_party_presolve.cpp b/cpp/src/mip/presolve/third_party_presolve.cpp index dc2d4b00e..88de4912a 100644 --- a/cpp/src/mip/presolve/third_party_presolve.cpp +++ b/cpp/src/mip/presolve/third_party_presolve.cpp @@ -274,19 +274,19 @@ void check_presolve_status(const papilo::PresolveStatus& status) { switch (status) { case papilo::PresolveStatus::kUnchanged: - CUOPT_LOG_INFO("Presolve status:: did not result in any changes"); + CUOPT_LOG_INFO("Presolve status: did not result in any changes"); break; case papilo::PresolveStatus::kReduced: - CUOPT_LOG_INFO("Presolve status:: reduced the problem"); + CUOPT_LOG_INFO("Presolve status: reduced the problem"); break; case papilo::PresolveStatus::kUnbndOrInfeas: - CUOPT_LOG_INFO("Presolve status:: found an unbounded or infeasible problem"); + CUOPT_LOG_INFO("Presolve status: found an unbounded or infeasible problem"); break; case papilo::PresolveStatus::kInfeasible: - CUOPT_LOG_INFO("Presolve status:: found an infeasible problem"); + CUOPT_LOG_INFO("Presolve status: found an infeasible problem"); break; case papilo::PresolveStatus::kUnbounded: - CUOPT_LOG_INFO("Presolve status:: found an unbounded problem"); + CUOPT_LOG_INFO("Presolve status: found an unbounded problem"); break; } } @@ -294,10 +294,10 @@ void check_presolve_status(const papilo::PresolveStatus& status) void check_postsolve_status(const papilo::PostsolveStatus& status) { switch (status) { - case papilo::PostsolveStatus::kOk: CUOPT_LOG_INFO("Post-solve status:: succeeded"); break; + case papilo::PostsolveStatus::kOk: CUOPT_LOG_INFO("Post-solve status: succeeded"); break; case papilo::PostsolveStatus::kFailed: CUOPT_LOG_INFO( - "Post-solve status:: Post solved solution violates constraints. This is most likely due to " + "Post-solve status: Post solved solution violates constraints. This is most likely due to " "different tolerances."); break; } @@ -362,7 +362,7 @@ std::pair, bool> third_party_presolve_t papilo_problem = build_papilo_problem(op_problem); - CUOPT_LOG_INFO("Unpresolved problem:: %d constraints, %d variables, %d nonzeros", + CUOPT_LOG_INFO("Unpresolved problem: %d constraints, %d variables, %d nonzeros", papilo_problem.getNRows(), papilo_problem.getNCols(), papilo_problem.getConstraintMatrix().getNnz()); @@ -382,11 +382,11 @@ std::pair, bool> third_party_presolve_t(op_problem.get_handle_ptr()), false); } post_solve_storage_ = result.postsolve; - CUOPT_LOG_INFO("Presolve removed:: %d constraints, %d variables, %d nonzeros", + CUOPT_LOG_INFO("Presolve removed: %d constraints, %d variables, %d nonzeros", op_problem.get_n_constraints() - papilo_problem.getNRows(), op_problem.get_n_variables() - papilo_problem.getNCols(), op_problem.get_nnz() - papilo_problem.getConstraintMatrix().getNnz()); - CUOPT_LOG_INFO("Presolved problem:: %d constraints, %d variables, %d nonzeros", + CUOPT_LOG_INFO("Presolved problem: %d constraints, %d variables, %d nonzeros", papilo_problem.getNRows(), papilo_problem.getNCols(), papilo_problem.getConstraintMatrix().getNnz());