diff --git a/ci/check_symbols.sh b/ci/check_symbols.sh new file mode 100755 index 0000000000..e35ac308f0 --- /dev/null +++ b/ci/check_symbols.sh @@ -0,0 +1,44 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. +# SPDX-License-Identifier: Apache-2.0 + +set -eEuo pipefail + +echo "checking for symbol visibility issues" + +LIBRARY="${1}" + +echo "" +echo "Checking exported symbols in '${LIBRARY}'" +symbol_file="$(mktemp)" +readelf --dyn-syms --wide "${LIBRARY}" \ + | awk '$7 != "UND"' \ + | c++filt \ + > "${symbol_file}" + +patterns=( + 'cub::' + 'thrust::' + 'raft::' + 'rmm::' + 'cuopt::linear_programming::detail' + 'cuopt::routing::detail' + 'grpc::' + 'google::protobuf' + 'tbb::' +) + +for pattern in "${patterns[@]}"; do + echo "Checking for '${pattern}' symbols..." + matches=$(grep -F -c "${pattern}" "${symbol_file}" || true) + if [[ "${matches}" -ne 0 ]]; then + grep -F -m 20 "${pattern}" "${symbol_file}" + echo "ERROR: Found exported symbols in ${LIBRARY} matching the pattern ${pattern}." + echo "ERROR: Total matching symbols: ${matches}" + rm "${symbol_file}" + exit 1 + fi +done + +rm "${symbol_file}" +echo "No symbol visibility issues found in ${LIBRARY}" diff --git a/conda/recipes/libcuopt/recipe.yaml b/conda/recipes/libcuopt/recipe.yaml index ee074392ae..4bf1b9120d 100644 --- a/conda/recipes/libcuopt/recipe.yaml +++ b/conda/recipes/libcuopt/recipe.yaml @@ -106,6 +106,7 @@ outputs: script: content: | cmake --install cpp/libmps_parser/build + ./ci/check_symbols.sh "$PREFIX/lib/libmps_parser.so" dynamic_linking: overlinking_behavior: "error" prefix_detection: @@ -159,6 +160,7 @@ outputs: script: content: | cmake --install cpp/build + ./ci/check_symbols.sh "$PREFIX/lib/libcuopt.so" dynamic_linking: overlinking_behavior: "error" prefix_detection: diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 57637bf1e6..bfc701f5bf 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -185,9 +185,6 @@ message("-- Host target architecture = '${CMAKE_SYSTEM_PROCESSOR}'") # make the flags global in order to propagate flags to test cmake files set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr --expt-extended-lambda") -if (${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.9) - set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -static-global-template-stub=false") -endif () list(APPEND CUOPT_CUDA_FLAGS -Werror=cross-execution-space-call -Wno-deprecated-declarations -Xcompiler=-Werror --default-stream=per-thread) if ("${CMAKE_CUDA_HOST_COMPILER}" MATCHES "clang" OR "${CMAKE_CXX_COMPILER_ID}" STREQUAL "Clang") list(APPEND CUOPT_CUDA_FLAGS -Xcompiler=-Wall) @@ -435,6 +432,8 @@ if (NOT SKIP_GRPC_BUILD) # at runtime with "undefined symbol: absl::…::Mutex::Dtor". set_property(SOURCE ${GRPC_INFRA_FILES} DIRECTORY ${CMAKE_SOURCE_DIR} APPEND PROPERTY COMPILE_OPTIONS "-DNDEBUG") + set_property(SOURCE ${PROTO_SRCS} ${GRPC_PROTO_SRCS} ${GRPC_SERVICE_SRCS} DIRECTORY ${CMAKE_SOURCE_DIR} + APPEND PROPERTY COMPILE_OPTIONS "$<$:-fvisibility=default>") endif (NOT SKIP_GRPC_BUILD) add_library(cuopt SHARED @@ -448,6 +447,14 @@ set_target_properties(cuopt CXX_SCAN_FOR_MODULES OFF ) +if (NOT BUILD_TESTS) + set_target_properties(cuopt + PROPERTIES CXX_VISIBILITY_PRESET hidden + CUDA_VISIBILITY_PRESET hidden + VISIBILITY_INLINES_HIDDEN ON + ) +endif () + target_compile_definitions(cuopt PUBLIC "CUOPT_LOG_ACTIVE_LEVEL=RAPIDS_LOGGER_LOG_LEVEL_${LIBCUOPT_LOGGING_LEVEL}") target_compile_options(cuopt diff --git a/cpp/include/cuopt/common/export.hpp b/cpp/include/cuopt/common/export.hpp new file mode 100644 index 0000000000..5f7209c5f5 --- /dev/null +++ b/cpp/include/cuopt/common/export.hpp @@ -0,0 +1,14 @@ +/* clang-format off */ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +/* clang-format on */ + +#pragma once + +#if defined(__GNUC__) || defined(__clang__) +#define CUOPT_EXPORT __attribute__((visibility("default"))) +#else +#define CUOPT_EXPORT +#endif diff --git a/cpp/include/cuopt/linear_programming/backend_selection.hpp b/cpp/include/cuopt/linear_programming/backend_selection.hpp index 26bcd942a3..d5addabca7 100644 --- a/cpp/include/cuopt/linear_programming/backend_selection.hpp +++ b/cpp/include/cuopt/linear_programming/backend_selection.hpp @@ -7,7 +7,10 @@ #pragma once -namespace cuopt::linear_programming { +#include + +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { /** * @brief Enum for execution mode (local vs remote solve) @@ -61,4 +64,5 @@ bool use_cpu_memory_for_local(); */ memory_backend_t get_memory_backend_type(); -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/cpu_optimization_problem.hpp b/cpp/include/cuopt/linear_programming/cpu_optimization_problem.hpp index 48d61b9e0c..7e7ea7540e 100644 --- a/cpp/include/cuopt/linear_programming/cpu_optimization_problem.hpp +++ b/cpp/include/cuopt/linear_programming/cpu_optimization_problem.hpp @@ -7,6 +7,7 @@ #pragma once +#include #include #include @@ -17,7 +18,8 @@ #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { // Forward declarations template @@ -206,4 +208,5 @@ class cpu_optimization_problem_t : public optimization_problem_interface_t row_names_{}; }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/cpu_optimization_problem_solution.hpp b/cpp/include/cuopt/linear_programming/cpu_optimization_problem_solution.hpp index e86dd0341a..42cc21cce5 100644 --- a/cpp/include/cuopt/linear_programming/cpu_optimization_problem_solution.hpp +++ b/cpp/include/cuopt/linear_programming/cpu_optimization_problem_solution.hpp @@ -7,6 +7,7 @@ #pragma once +#include #include #include #include @@ -18,7 +19,8 @@ #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { /** * @brief CPU-backed LP solution (uses std::vector instead of rmm::device_uvector) @@ -389,4 +391,5 @@ class cpu_mip_solution_t : public mip_solution_interface_t { i_t num_simplex_iterations_; }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/cpu_pdlp_warm_start_data.hpp b/cpp/include/cuopt/linear_programming/cpu_pdlp_warm_start_data.hpp index 51e98275ed..952e779fdc 100644 --- a/cpp/include/cuopt/linear_programming/cpu_pdlp_warm_start_data.hpp +++ b/cpp/include/cuopt/linear_programming/cpu_pdlp_warm_start_data.hpp @@ -7,10 +7,12 @@ #pragma once +#include #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { // CPU version of pdlp_warm_start_data_t using std::vector for remote execution template @@ -118,4 +120,5 @@ template pdlp_warm_start_data_t convert_to_gpu_warmstart( const cpu_pdlp_warm_start_data_t& cpu_data, rmm::cuda_stream_view stream); -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/cuopt_c.h b/cpp/include/cuopt/linear_programming/cuopt_c.h index 4c4d44c764..af9fe7141a 100644 --- a/cpp/include/cuopt/linear_programming/cuopt_c.h +++ b/cpp/include/cuopt/linear_programming/cuopt_c.h @@ -9,6 +9,7 @@ #define CUOPT_C_API_H #include +#include #include @@ -76,12 +77,12 @@ typedef int64_t cuopt_int_t; * * @return The size in bytes of the float type. */ -int8_t cuOptGetFloatSize(); +CUOPT_EXPORT int8_t cuOptGetFloatSize(); /** @brief Get the size of the integer type used by the library. * @return The size of the integer type in bytes. */ -int8_t cuOptGetIntSize(); +CUOPT_EXPORT int8_t cuOptGetIntSize(); /** * @brief Get the version of the library. @@ -95,9 +96,9 @@ int8_t cuOptGetIntSize(); * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetVersion(cuopt_int_t* version_major, - cuopt_int_t* version_minor, - cuopt_int_t* version_patch); +CUOPT_EXPORT cuopt_int_t cuOptGetVersion(cuopt_int_t* version_major, + cuopt_int_t* version_minor, + cuopt_int_t* version_patch); /** * @brief Read an optimization problem from an MPS file. @@ -109,7 +110,8 @@ cuopt_int_t cuOptGetVersion(cuopt_int_t* version_major, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptReadProblem(const char* filename, cuOptOptimizationProblem* problem_ptr); +CUOPT_EXPORT cuopt_int_t cuOptReadProblem(const char* filename, + cuOptOptimizationProblem* problem_ptr); /** * @brief Write an optimization problem to a file. @@ -121,9 +123,9 @@ cuopt_int_t cuOptReadProblem(const char* filename, cuOptOptimizationProblem* pro * @return A status code indicating success or failure. Returns CUOPT_INVALID_ARGUMENT * if an unsupported format is specified. */ -cuopt_int_t cuOptWriteProblem(cuOptOptimizationProblem problem, - const char* filename, - cuopt_int_t format); +CUOPT_EXPORT cuopt_int_t cuOptWriteProblem(cuOptOptimizationProblem problem, + const char* filename, + cuopt_int_t format); /** @brief Create an optimization problem of the form * @@ -168,20 +170,21 @@ cuopt_int_t cuOptWriteProblem(cuOptOptimizationProblem problem, * @param[out] problem_ptr Pointer to store the created optimization problem * @return CUOPT_SUCCESS if successful, CUOPT_ERROR otherwise */ -cuopt_int_t cuOptCreateProblem(cuopt_int_t num_constraints, - cuopt_int_t num_variables, - cuopt_int_t objective_sense, - cuopt_float_t objective_offset, - const cuopt_float_t* objective_coefficients, - const cuopt_int_t* constraint_matrix_row_offsets, - const cuopt_int_t* constraint_matrix_column_indices, - const cuopt_float_t* constraint_matrix_coefficent_values, - const char* constraint_sense, - const cuopt_float_t* rhs, - const cuopt_float_t* lower_bounds, - const cuopt_float_t* upper_bounds, - const char* variable_types, - cuOptOptimizationProblem* problem_ptr); +CUOPT_EXPORT cuopt_int_t +cuOptCreateProblem(cuopt_int_t num_constraints, + cuopt_int_t num_variables, + cuopt_int_t objective_sense, + cuopt_float_t objective_offset, + const cuopt_float_t* objective_coefficients, + const cuopt_int_t* constraint_matrix_row_offsets, + const cuopt_int_t* constraint_matrix_column_indices, + const cuopt_float_t* constraint_matrix_coefficent_values, + const char* constraint_sense, + const cuopt_float_t* rhs, + const cuopt_float_t* lower_bounds, + const cuopt_float_t* upper_bounds, + const char* variable_types, + cuOptOptimizationProblem* problem_ptr); /** @brief Create an optimization problem of the form * * @verbatim @@ -237,20 +240,21 @@ cuopt_int_t cuOptCreateProblem(cuopt_int_t num_constraints, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptCreateRangedProblem(cuopt_int_t num_constraints, - cuopt_int_t num_variables, - cuopt_int_t objective_sense, - cuopt_float_t objective_offset, - const cuopt_float_t* objective_coefficients, - const cuopt_int_t* constraint_matrix_row_offsets, - const cuopt_int_t* constraint_matrix_column_indices, - const cuopt_float_t* constraint_matrix_coefficients, - const cuopt_float_t* constraint_lower_bounds, - const cuopt_float_t* constraint_upper_bounds, - const cuopt_float_t* variable_lower_bounds, - const cuopt_float_t* variable_upper_bounds, - const char* variable_types, - cuOptOptimizationProblem* problem_ptr); +CUOPT_EXPORT cuopt_int_t +cuOptCreateRangedProblem(cuopt_int_t num_constraints, + cuopt_int_t num_variables, + cuopt_int_t objective_sense, + cuopt_float_t objective_offset, + const cuopt_float_t* objective_coefficients, + const cuopt_int_t* constraint_matrix_row_offsets, + const cuopt_int_t* constraint_matrix_column_indices, + const cuopt_float_t* constraint_matrix_coefficients, + const cuopt_float_t* constraint_lower_bounds, + const cuopt_float_t* constraint_upper_bounds, + const cuopt_float_t* variable_lower_bounds, + const cuopt_float_t* variable_upper_bounds, + const char* variable_types, + cuOptOptimizationProblem* problem_ptr); /** @brief Create an optimization problem of the form * @@ -306,23 +310,23 @@ cuopt_int_t cuOptCreateRangedProblem(cuopt_int_t num_constraints, * @param[out] problem_ptr Pointer to store the created optimization problem * @return CUOPT_SUCCESS if successful, CUOPT_ERROR otherwise */ -cuopt_int_t cuOptCreateQuadraticProblem( - cuopt_int_t num_constraints, - cuopt_int_t num_variables, - cuopt_int_t objective_sense, - cuopt_float_t objective_offset, - const cuopt_float_t* objective_coefficients, - const cuopt_int_t* quadratic_objective_matrix_row_offsets, - const cuopt_int_t* quadratic_objective_matrix_column_indices, - const cuopt_float_t* quadratic_objective_matrix_coefficent_values, - const cuopt_int_t* constraint_matrix_row_offsets, - const cuopt_int_t* constraint_matrix_column_indices, - const cuopt_float_t* constraint_matrix_coefficent_values, - const char* constraint_sense, - const cuopt_float_t* rhs, - const cuopt_float_t* lower_bounds, - const cuopt_float_t* upper_bounds, - cuOptOptimizationProblem* problem_ptr); +CUOPT_EXPORT cuopt_int_t +cuOptCreateQuadraticProblem(cuopt_int_t num_constraints, + cuopt_int_t num_variables, + cuopt_int_t objective_sense, + cuopt_float_t objective_offset, + const cuopt_float_t* objective_coefficients, + const cuopt_int_t* quadratic_objective_matrix_row_offsets, + const cuopt_int_t* quadratic_objective_matrix_column_indices, + const cuopt_float_t* quadratic_objective_matrix_coefficent_values, + const cuopt_int_t* constraint_matrix_row_offsets, + const cuopt_int_t* constraint_matrix_column_indices, + const cuopt_float_t* constraint_matrix_coefficent_values, + const char* constraint_sense, + const cuopt_float_t* rhs, + const cuopt_float_t* lower_bounds, + const cuopt_float_t* upper_bounds, + cuOptOptimizationProblem* problem_ptr); /** @brief Create an optimization problem of the form * * @verbatim @@ -388,30 +392,30 @@ cuopt_int_t cuOptCreateQuadraticProblem( * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptCreateQuadraticRangedProblem( - cuopt_int_t num_constraints, - cuopt_int_t num_variables, - cuopt_int_t objective_sense, - cuopt_float_t objective_offset, - const cuopt_float_t* objective_coefficients, - const cuopt_int_t* quadratic_objective_matrix_row_offsets, - const cuopt_int_t* quadratic_objective_matrix_column_indices, - const cuopt_float_t* quadratic_objective_matrix_coefficent_values, - const cuopt_int_t* constraint_matrix_row_offsets, - const cuopt_int_t* constraint_matrix_column_indices, - const cuopt_float_t* constraint_matrix_coefficients, - const cuopt_float_t* constraint_lower_bounds, - const cuopt_float_t* constraint_upper_bounds, - const cuopt_float_t* variable_lower_bounds, - const cuopt_float_t* variable_upper_bounds, - cuOptOptimizationProblem* problem_ptr); +CUOPT_EXPORT cuopt_int_t +cuOptCreateQuadraticRangedProblem(cuopt_int_t num_constraints, + cuopt_int_t num_variables, + cuopt_int_t objective_sense, + cuopt_float_t objective_offset, + const cuopt_float_t* objective_coefficients, + const cuopt_int_t* quadratic_objective_matrix_row_offsets, + const cuopt_int_t* quadratic_objective_matrix_column_indices, + const cuopt_float_t* quadratic_objective_matrix_coefficent_values, + const cuopt_int_t* constraint_matrix_row_offsets, + const cuopt_int_t* constraint_matrix_column_indices, + const cuopt_float_t* constraint_matrix_coefficients, + const cuopt_float_t* constraint_lower_bounds, + const cuopt_float_t* constraint_upper_bounds, + const cuopt_float_t* variable_lower_bounds, + const cuopt_float_t* variable_upper_bounds, + cuOptOptimizationProblem* problem_ptr); /** @brief Destroy an optimization problem * * @param[in, out] problem_ptr - A pointer to a cuOptOptimizationProblem. On * output the problem will be destroyed, and the pointer will be set to NULL. */ -void cuOptDestroyProblem(cuOptOptimizationProblem* problem_ptr); +CUOPT_EXPORT void cuOptDestroyProblem(cuOptOptimizationProblem* problem_ptr); /** @brief Get the number of constraints of an optimization problem. * @@ -422,8 +426,8 @@ void cuOptDestroyProblem(cuOptOptimizationProblem* problem_ptr); * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetNumConstraints(cuOptOptimizationProblem problem, - cuopt_int_t* num_constraints_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetNumConstraints(cuOptOptimizationProblem problem, + cuopt_int_t* num_constraints_ptr); /** @brief Get the number of variables of an optimization problem. * @@ -434,7 +438,8 @@ cuopt_int_t cuOptGetNumConstraints(cuOptOptimizationProblem problem, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetNumVariables(cuOptOptimizationProblem problem, cuopt_int_t* num_variables_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetNumVariables(cuOptOptimizationProblem problem, + cuopt_int_t* num_variables_ptr); /** @brief Get the objective sense of an optimization problem. * @@ -445,8 +450,8 @@ cuopt_int_t cuOptGetNumVariables(cuOptOptimizationProblem problem, cuopt_int_t* * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetObjectiveSense(cuOptOptimizationProblem problem, - cuopt_int_t* objective_sense_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetObjectiveSense(cuOptOptimizationProblem problem, + cuopt_int_t* objective_sense_ptr); /** @brief Get the objective offset of an optimization problem. * @@ -457,8 +462,8 @@ cuopt_int_t cuOptGetObjectiveSense(cuOptOptimizationProblem problem, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetObjectiveOffset(cuOptOptimizationProblem problem, - cuopt_float_t* objective_offset_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetObjectiveOffset(cuOptOptimizationProblem problem, + cuopt_float_t* objective_offset_ptr); /** @brief Get the objective coefficients of an optimization problem. * @@ -470,8 +475,8 @@ cuopt_int_t cuOptGetObjectiveOffset(cuOptOptimizationProblem problem, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetObjectiveCoefficients(cuOptOptimizationProblem problem, - cuopt_float_t* objective_coefficients_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetObjectiveCoefficients(cuOptOptimizationProblem problem, + cuopt_float_t* objective_coefficients_ptr); /** @brief Get the number of non-zero elements in the constraint matrix of an * optimization problem. @@ -483,7 +488,8 @@ cuopt_int_t cuOptGetObjectiveCoefficients(cuOptOptimizationProblem problem, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetNumNonZeros(cuOptOptimizationProblem problem, cuopt_int_t* num_non_zeros_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetNumNonZeros(cuOptOptimizationProblem problem, + cuopt_int_t* num_non_zeros_ptr); /** @brief Get the constraint matrix of an optimization problem in compressed sparse row format. * @@ -503,10 +509,11 @@ cuopt_int_t cuOptGetNumNonZeros(cuOptOptimizationProblem problem, cuopt_int_t* n * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetConstraintMatrix(cuOptOptimizationProblem problem, - cuopt_int_t* constraint_matrix_row_offsets_ptr, - cuopt_int_t* constraint_matrix_column_indices_ptr, - cuopt_float_t* constraint_matrix_coefficients_ptr); +CUOPT_EXPORT cuopt_int_t +cuOptGetConstraintMatrix(cuOptOptimizationProblem problem, + cuopt_int_t* constraint_matrix_row_offsets_ptr, + cuopt_int_t* constraint_matrix_column_indices_ptr, + cuopt_float_t* constraint_matrix_coefficients_ptr); /** @brief Get the constraint sense of an optimization problem. * @@ -517,7 +524,8 @@ cuopt_int_t cuOptGetConstraintMatrix(cuOptOptimizationProblem problem, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetConstraintSense(cuOptOptimizationProblem problem, char* constraint_sense_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetConstraintSense(cuOptOptimizationProblem problem, + char* constraint_sense_ptr); /** @brief Get the right-hand side of an optimization problem. * @@ -528,8 +536,8 @@ cuopt_int_t cuOptGetConstraintSense(cuOptOptimizationProblem problem, char* cons * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetConstraintRightHandSide(cuOptOptimizationProblem problem, - cuopt_float_t* rhs_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetConstraintRightHandSide(cuOptOptimizationProblem problem, + cuopt_float_t* rhs_ptr); /** @brief Get the lower bounds of an optimization problem. * @@ -540,8 +548,8 @@ cuopt_int_t cuOptGetConstraintRightHandSide(cuOptOptimizationProblem problem, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetConstraintLowerBounds(cuOptOptimizationProblem problem, - cuopt_float_t* lower_bounds_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetConstraintLowerBounds(cuOptOptimizationProblem problem, + cuopt_float_t* lower_bounds_ptr); /** @brief Get the upper bounds of an optimization problem. * @@ -552,8 +560,8 @@ cuopt_int_t cuOptGetConstraintLowerBounds(cuOptOptimizationProblem problem, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetConstraintUpperBounds(cuOptOptimizationProblem problem, - cuopt_float_t* upper_bounds_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetConstraintUpperBounds(cuOptOptimizationProblem problem, + cuopt_float_t* upper_bounds_ptr); /** @brief Get the lower bounds of an optimization problem. * @@ -564,8 +572,8 @@ cuopt_int_t cuOptGetConstraintUpperBounds(cuOptOptimizationProblem problem, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetVariableLowerBounds(cuOptOptimizationProblem problem, - cuopt_float_t* lower_bounds_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetVariableLowerBounds(cuOptOptimizationProblem problem, + cuopt_float_t* lower_bounds_ptr); /** @brief Get the upper bounds of an optimization problem. * @@ -576,8 +584,8 @@ cuopt_int_t cuOptGetVariableLowerBounds(cuOptOptimizationProblem problem, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetVariableUpperBounds(cuOptOptimizationProblem problem, - cuopt_float_t* upper_bounds_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetVariableUpperBounds(cuOptOptimizationProblem problem, + cuopt_float_t* upper_bounds_ptr); /** @brief Get the variable types of an optimization problem. * @@ -589,7 +597,8 @@ cuopt_int_t cuOptGetVariableUpperBounds(cuOptOptimizationProblem problem, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetVariableTypes(cuOptOptimizationProblem problem, char* variable_types_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetVariableTypes(cuOptOptimizationProblem problem, + char* variable_types_ptr); /** @brief Create a solver settings object. * @@ -598,14 +607,14 @@ cuopt_int_t cuOptGetVariableTypes(cuOptOptimizationProblem problem, char* variab * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptCreateSolverSettings(cuOptSolverSettings* settings_ptr); +CUOPT_EXPORT cuopt_int_t cuOptCreateSolverSettings(cuOptSolverSettings* settings_ptr); /** @brief Destroy a solver settings object. * * @param[in, out] settings_ptr - A pointer to a cuOptSolverSettings object. On output * the solver settings will be destroyed and the pointer will be set to NULL. */ -void cuOptDestroySolverSettings(cuOptSolverSettings* settings_ptr); +CUOPT_EXPORT void cuOptDestroySolverSettings(cuOptSolverSettings* settings_ptr); /** @brief Set a parameter of a solver settings object. * @@ -615,9 +624,9 @@ void cuOptDestroySolverSettings(cuOptSolverSettings* settings_ptr); * * @param[in] parameter_value - The value of the parameter to set. */ -cuopt_int_t cuOptSetParameter(cuOptSolverSettings settings, - const char* parameter_name, - const char* parameter_value); +CUOPT_EXPORT cuopt_int_t cuOptSetParameter(cuOptSolverSettings settings, + const char* parameter_name, + const char* parameter_value); /** @brief Get a parameter of a solver settings object. * @@ -632,10 +641,10 @@ cuopt_int_t cuOptSetParameter(cuOptSolverSettings settings, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetParameter(cuOptSolverSettings settings, - const char* parameter_name, - cuopt_int_t parameter_value_size, - char* parameter_value); +CUOPT_EXPORT cuopt_int_t cuOptGetParameter(cuOptSolverSettings settings, + const char* parameter_name, + cuopt_int_t parameter_value_size, + char* parameter_value); /** @brief Set an integer parameter of a solver settings object. * @@ -647,9 +656,9 @@ cuopt_int_t cuOptGetParameter(cuOptSolverSettings settings, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptSetIntegerParameter(cuOptSolverSettings settings, - const char* parameter_name, - cuopt_int_t parameter_value); +CUOPT_EXPORT cuopt_int_t cuOptSetIntegerParameter(cuOptSolverSettings settings, + const char* parameter_name, + cuopt_int_t parameter_value); /** @brief Get an integer parameter of a solver settings object. * @@ -662,9 +671,9 @@ cuopt_int_t cuOptSetIntegerParameter(cuOptSolverSettings settings, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetIntegerParameter(cuOptSolverSettings settings, - const char* parameter_name, - cuopt_int_t* parameter_value); +CUOPT_EXPORT cuopt_int_t cuOptGetIntegerParameter(cuOptSolverSettings settings, + const char* parameter_name, + cuopt_int_t* parameter_value); /** @brief Set a float parameter of a solver settings object. * @@ -676,9 +685,9 @@ cuopt_int_t cuOptGetIntegerParameter(cuOptSolverSettings settings, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptSetFloatParameter(cuOptSolverSettings settings, - const char* parameter_name, - cuopt_float_t parameter_value); +CUOPT_EXPORT cuopt_int_t cuOptSetFloatParameter(cuOptSolverSettings settings, + const char* parameter_name, + cuopt_float_t parameter_value); /** @brief Get a float parameter of a solver settings object. * @@ -691,9 +700,9 @@ cuopt_int_t cuOptSetFloatParameter(cuOptSolverSettings settings, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetFloatParameter(cuOptSolverSettings settings, - const char* parameter_name, - cuopt_float_t* parameter_value); +CUOPT_EXPORT cuopt_int_t cuOptGetFloatParameter(cuOptSolverSettings settings, + const char* parameter_name, + cuopt_float_t* parameter_value); /** * @brief Type of callback for receiving incumbent MIP solutions with user context. @@ -744,9 +753,9 @@ typedef void (*cuOptMIPSetSolutionCallback)(cuopt_float_t* solution, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptSetMIPGetSolutionCallback(cuOptSolverSettings settings, - cuOptMIPGetSolutionCallback callback, - void* user_data); +CUOPT_EXPORT cuopt_int_t cuOptSetMIPGetSolutionCallback(cuOptSolverSettings settings, + cuOptMIPGetSolutionCallback callback, + void* user_data); /** * @brief Register a callback to inject MIP solutions. @@ -762,9 +771,9 @@ cuopt_int_t cuOptSetMIPGetSolutionCallback(cuOptSolverSettings settings, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptSetMIPSetSolutionCallback(cuOptSolverSettings settings, - cuOptMIPSetSolutionCallback callback, - void* user_data); +CUOPT_EXPORT cuopt_int_t cuOptSetMIPSetSolutionCallback(cuOptSolverSettings settings, + cuOptMIPSetSolutionCallback callback, + void* user_data); /** * @brief Set the initial primal solution for an LP solve. * @@ -778,9 +787,9 @@ cuopt_int_t cuOptSetMIPSetSolutionCallback(cuOptSolverSettings settings, * @note All pointer arguments (primal_solution) refer to host memory. * @return A status code indicating success or failure. */ -cuopt_int_t cuOptSetInitialPrimalSolution(cuOptSolverSettings settings, - const cuopt_float_t* primal_solution, - cuopt_int_t num_variables); +CUOPT_EXPORT cuopt_int_t cuOptSetInitialPrimalSolution(cuOptSolverSettings settings, + const cuopt_float_t* primal_solution, + cuopt_int_t num_variables); /** * @brief Set the initial dual solution for an LP solve. @@ -795,9 +804,9 @@ cuopt_int_t cuOptSetInitialPrimalSolution(cuOptSolverSettings settings, * @note All pointer arguments (dual_solution) refer to host memory. * @return A status code indicating success or failure. */ -cuopt_int_t cuOptSetInitialDualSolution(cuOptSolverSettings settings, - const cuopt_float_t* dual_solution, - cuopt_int_t num_constraints); +CUOPT_EXPORT cuopt_int_t cuOptSetInitialDualSolution(cuOptSolverSettings settings, + const cuopt_float_t* dual_solution, + cuopt_int_t num_constraints); /** * @brief Add an initial solution (MIP start) for MIP solving. @@ -815,9 +824,9 @@ cuopt_int_t cuOptSetInitialDualSolution(cuOptSolverSettings settings, * @note All pointer arguments (solution) refer to host memory. * @return A status code indicating success or failure. */ -cuopt_int_t cuOptAddMIPStart(cuOptSolverSettings settings, - const cuopt_float_t* solution, - cuopt_int_t num_variables); +CUOPT_EXPORT cuopt_int_t cuOptAddMIPStart(cuOptSolverSettings settings, + const cuopt_float_t* solution, + cuopt_int_t num_variables); /** @brief Check if an optimization problem is a mixed integer programming problem. * @@ -828,7 +837,7 @@ cuopt_int_t cuOptAddMIPStart(cuOptSolverSettings settings, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptIsMIP(cuOptOptimizationProblem problem, cuopt_int_t* is_mip_ptr); +CUOPT_EXPORT cuopt_int_t cuOptIsMIP(cuOptOptimizationProblem problem, cuopt_int_t* is_mip_ptr); /** @brief Solve an optimization problem. * @@ -841,16 +850,16 @@ cuopt_int_t cuOptIsMIP(cuOptOptimizationProblem problem, cuopt_int_t* is_mip_ptr * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptSolve(cuOptOptimizationProblem problem, - cuOptSolverSettings settings, - cuOptSolution* solution_ptr); +CUOPT_EXPORT cuopt_int_t cuOptSolve(cuOptOptimizationProblem problem, + cuOptSolverSettings settings, + cuOptSolution* solution_ptr); /** @brief Destroy a solution object. * * @param[in, out] solution_ptr - A pointer to a cuOptSolution object. On output * the solution will be destroyed and the pointer will be set to NULL. */ -void cuOptDestroySolution(cuOptSolution* solution_ptr); +CUOPT_EXPORT void cuOptDestroySolution(cuOptSolution* solution_ptr); /** @brief Get the termination reason of an optimization problem. * @@ -861,7 +870,8 @@ void cuOptDestroySolution(cuOptSolution* solution_ptr); * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetTerminationStatus(cuOptSolution solution, cuopt_int_t* termination_status_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetTerminationStatus(cuOptSolution solution, + cuopt_int_t* termination_status_ptr); /* @brief Get the error status of a solution object. * @@ -872,7 +882,7 @@ cuopt_int_t cuOptGetTerminationStatus(cuOptSolution solution, cuopt_int_t* termi * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetErrorStatus(cuOptSolution solution, cuopt_int_t* error_status_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetErrorStatus(cuOptSolution solution, cuopt_int_t* error_status_ptr); /* @brief Get the error string of a solution object. * @@ -885,9 +895,9 @@ cuopt_int_t cuOptGetErrorStatus(cuOptSolution solution, cuopt_int_t* error_statu * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetErrorString(cuOptSolution solution, - char* error_string_ptr, - cuopt_int_t error_string_size); +CUOPT_EXPORT cuopt_int_t cuOptGetErrorString(cuOptSolution solution, + char* error_string_ptr, + cuopt_int_t error_string_size); /* @brief Get the solution of an optimization problem. * @@ -898,7 +908,8 @@ cuopt_int_t cuOptGetErrorString(cuOptSolution solution, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetPrimalSolution(cuOptSolution solution, cuopt_float_t* solution_values); +CUOPT_EXPORT cuopt_int_t cuOptGetPrimalSolution(cuOptSolution solution, + cuopt_float_t* solution_values); /** @brief Get the objective value of an optimization problem. * @@ -909,7 +920,8 @@ cuopt_int_t cuOptGetPrimalSolution(cuOptSolution solution, cuopt_float_t* soluti * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetObjectiveValue(cuOptSolution solution, cuopt_float_t* objective_value_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetObjectiveValue(cuOptSolution solution, + cuopt_float_t* objective_value_ptr); /** @brief Get the solve time of an optimization problem. * @@ -919,7 +931,7 @@ cuopt_int_t cuOptGetObjectiveValue(cuOptSolution solution, cuopt_float_t* object * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetSolveTime(cuOptSolution solution, cuopt_float_t* solve_time_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetSolveTime(cuOptSolution solution, cuopt_float_t* solve_time_ptr); /** @brief Get the relative MIP gap of an optimization problem. * @@ -930,7 +942,7 @@ cuopt_int_t cuOptGetSolveTime(cuOptSolution solution, cuopt_float_t* solve_time_ * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetMIPGap(cuOptSolution solution, cuopt_float_t* mip_gap_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetMIPGap(cuOptSolution solution, cuopt_float_t* mip_gap_ptr); /** @brief Get the solution bound of an optimization problem. * @@ -941,7 +953,8 @@ cuopt_int_t cuOptGetMIPGap(cuOptSolution solution, cuopt_float_t* mip_gap_ptr); * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetSolutionBound(cuOptSolution solution, cuopt_float_t* solution_bound_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetSolutionBound(cuOptSolution solution, + cuopt_float_t* solution_bound_ptr); /** @brief Get the dual solution of an optimization problem. * @@ -952,7 +965,8 @@ cuopt_int_t cuOptGetSolutionBound(cuOptSolution solution, cuopt_float_t* solutio * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetDualSolution(cuOptSolution solution, cuopt_float_t* dual_solution_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetDualSolution(cuOptSolution solution, + cuopt_float_t* dual_solution_ptr); /** @brief Get the dual objective value of an optimization problem. * @@ -963,8 +977,8 @@ cuopt_int_t cuOptGetDualSolution(cuOptSolution solution, cuopt_float_t* dual_sol * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetDualObjectiveValue(cuOptSolution solution, - cuopt_float_t* dual_objective_value_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetDualObjectiveValue(cuOptSolution solution, + cuopt_float_t* dual_objective_value_ptr); /** @brief Get the reduced costs of an optimization problem. * @@ -975,7 +989,8 @@ cuopt_int_t cuOptGetDualObjectiveValue(cuOptSolution solution, * * @return A status code indicating success or failure. */ -cuopt_int_t cuOptGetReducedCosts(cuOptSolution solution, cuopt_float_t* reduced_cost_ptr); +CUOPT_EXPORT cuopt_int_t cuOptGetReducedCosts(cuOptSolution solution, + cuopt_float_t* reduced_cost_ptr); #ifdef __cplusplus } diff --git a/cpp/include/cuopt/linear_programming/mip/solver_settings.hpp b/cpp/include/cuopt/linear_programming/mip/solver_settings.hpp index ae0187e454..ae7a8a77a7 100644 --- a/cpp/include/cuopt/linear_programming/mip/solver_settings.hpp +++ b/cpp/include/cuopt/linear_programming/mip/solver_settings.hpp @@ -7,6 +7,8 @@ #pragma once +#include + #include #include @@ -20,7 +22,8 @@ #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { struct benchmark_info_t { double last_improvement_of_best_feasible = 0; @@ -200,4 +203,5 @@ struct mip_solver_settings_accessor { } // namespace detail -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/mip/solver_solution.hpp b/cpp/include/cuopt/linear_programming/mip/solver_solution.hpp index b8fa884540..68d96d5b65 100644 --- a/cpp/include/cuopt/linear_programming/mip/solver_solution.hpp +++ b/cpp/include/cuopt/linear_programming/mip/solver_solution.hpp @@ -8,6 +8,7 @@ #pragma once #include +#include #include #include #include @@ -21,7 +22,8 @@ #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { enum class mip_termination_status_t : int8_t { NoTermination = CUOPT_TERMINATION_STATUS_NO_TERMINATION, @@ -92,4 +94,5 @@ class mip_solution_t : public base_solution_t { std::vector> solution_pool_; }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/optimization_problem.hpp b/cpp/include/cuopt/linear_programming/optimization_problem.hpp index a61118aa1c..bcc8c406ac 100644 --- a/cpp/include/cuopt/linear_programming/optimization_problem.hpp +++ b/cpp/include/cuopt/linear_programming/optimization_problem.hpp @@ -7,6 +7,7 @@ #pragma once +#include #include #include @@ -19,7 +20,8 @@ #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { // Forward declarations template @@ -417,4 +419,5 @@ class optimization_problem_t : public optimization_problem_interface_t std::vector row_names_{}; }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/optimization_problem_solution.hpp b/cpp/include/cuopt/linear_programming/optimization_problem_solution.hpp index ac55256973..1d918e079a 100644 --- a/cpp/include/cuopt/linear_programming/optimization_problem_solution.hpp +++ b/cpp/include/cuopt/linear_programming/optimization_problem_solution.hpp @@ -7,6 +7,7 @@ #pragma once +#include #include #include #include @@ -16,7 +17,8 @@ #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { /** * @brief GPU-backed LP solution (wraps optimization_problem_solution_t) @@ -476,4 +478,5 @@ class gpu_mip_solution_t : public mip_solution_interface_t { mip_solution_t solution_; }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/pdlp/pdlp_warm_start_data.hpp b/cpp/include/cuopt/linear_programming/pdlp/pdlp_warm_start_data.hpp index 1f241463ac..7ca5d836b2 100644 --- a/cpp/include/cuopt/linear_programming/pdlp/pdlp_warm_start_data.hpp +++ b/cpp/include/cuopt/linear_programming/pdlp/pdlp_warm_start_data.hpp @@ -7,11 +7,14 @@ #pragma once +#include + #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { template struct pdlp_warm_start_data_view_t; @@ -99,4 +102,5 @@ struct pdlp_warm_start_data_view_t { i_t iterations_since_last_restart_{-1}; }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/pdlp/solver_settings.hpp b/cpp/include/cuopt/linear_programming/pdlp/solver_settings.hpp index a1cb787f09..0e60c76481 100644 --- a/cpp/include/cuopt/linear_programming/pdlp/solver_settings.hpp +++ b/cpp/include/cuopt/linear_programming/pdlp/solver_settings.hpp @@ -8,6 +8,7 @@ #pragma once #include +#include #include #include #include @@ -21,7 +22,8 @@ #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { // Forward declare solver_settings_t for friend class template @@ -349,4 +351,5 @@ class pdlp_solver_settings_t { friend class solver_settings_t; }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/pdlp/solver_solution.hpp b/cpp/include/cuopt/linear_programming/pdlp/solver_solution.hpp index 81e911df62..891bc6298c 100644 --- a/cpp/include/cuopt/linear_programming/pdlp/solver_solution.hpp +++ b/cpp/include/cuopt/linear_programming/pdlp/solver_solution.hpp @@ -8,6 +8,7 @@ #pragma once #include +#include #include #include #include @@ -22,7 +23,8 @@ #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { // Possible reasons for terminating enum class pdlp_termination_status_t : int8_t { @@ -310,4 +312,5 @@ class optimization_problem_solution_t : public base_solution_t { /** error struct */ cuopt::logic_error error_status_; }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/include/cuopt/linear_programming/solver_settings.hpp b/cpp/include/cuopt/linear_programming/solver_settings.hpp index 1720b0e9f9..bb232c3679 100644 --- a/cpp/include/cuopt/linear_programming/solver_settings.hpp +++ b/cpp/include/cuopt/linear_programming/solver_settings.hpp @@ -7,6 +7,7 @@ #pragma once +#include #include #include @@ -22,7 +23,8 @@ #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { template class solver_settings_t { @@ -109,4 +111,5 @@ class solver_settings_t { std::vector> string_parameters; }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/libmps_parser/CMakeLists.txt b/cpp/libmps_parser/CMakeLists.txt index cab37741ff..3617a77c1e 100644 --- a/cpp/libmps_parser/CMakeLists.txt +++ b/cpp/libmps_parser/CMakeLists.txt @@ -84,6 +84,7 @@ set_target_properties(mps_parser PROPERTIES BUILD_RPATH "\$ORIGIN" INSTALL_RPATH "\$ORIGIN" INTERFACE_POSITION_INDEPENDENT_CODE ON + CXX_VISIBILITY_PRESET hidden CXX_SCAN_FOR_MODULES OFF ) diff --git a/cpp/libmps_parser/include/mps_parser/data_model_view.hpp b/cpp/libmps_parser/include/mps_parser/data_model_view.hpp index 04ed4d6b7c..2b496941b2 100644 --- a/cpp/libmps_parser/include/mps_parser/data_model_view.hpp +++ b/cpp/libmps_parser/include/mps_parser/data_model_view.hpp @@ -7,6 +7,7 @@ #pragma once +#include #include #include @@ -15,7 +16,8 @@ #include #include -namespace cuopt::mps_parser { +namespace cuopt { +namespace MPS_PARSER_EXPORT mps_parser { /** * @brief A representation of a linear programming (LP) optimization problem @@ -478,4 +480,5 @@ class data_model_view_t { std::vector::quadratic_constraint_t> quadratic_constraints_; }; // class data_model_view_t -} // namespace cuopt::mps_parser +} // namespace MPS_PARSER_EXPORT mps_parser +} // namespace cuopt diff --git a/cpp/libmps_parser/include/mps_parser/export.hpp b/cpp/libmps_parser/include/mps_parser/export.hpp new file mode 100644 index 0000000000..51058e752c --- /dev/null +++ b/cpp/libmps_parser/include/mps_parser/export.hpp @@ -0,0 +1,14 @@ +/* clang-format off */ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +/* clang-format on */ + +#pragma once + +#if defined(__GNUC__) || defined(__clang__) +#define MPS_PARSER_EXPORT __attribute__((visibility("default"))) +#else +#define MPS_PARSER_EXPORT +#endif diff --git a/cpp/libmps_parser/include/mps_parser/mps_data_model.hpp b/cpp/libmps_parser/include/mps_parser/mps_data_model.hpp index 4ca56f02ba..de3e3fb890 100644 --- a/cpp/libmps_parser/include/mps_parser/mps_data_model.hpp +++ b/cpp/libmps_parser/include/mps_parser/mps_data_model.hpp @@ -7,13 +7,16 @@ #pragma once +#include + #include #include #include #include #include -namespace cuopt::mps_parser { +namespace cuopt { +namespace MPS_PARSER_EXPORT mps_parser { /** * @brief A representation of a linear programming (LP) optimization problem @@ -383,4 +386,5 @@ class mps_data_model_t { }; // class mps_data_model_t -} // namespace cuopt::mps_parser +} // namespace MPS_PARSER_EXPORT mps_parser +} // namespace cuopt diff --git a/cpp/libmps_parser/include/mps_parser/mps_writer.hpp b/cpp/libmps_parser/include/mps_parser/mps_writer.hpp index 30f2fdf942..aca3bde2d9 100644 --- a/cpp/libmps_parser/include/mps_parser/mps_writer.hpp +++ b/cpp/libmps_parser/include/mps_parser/mps_writer.hpp @@ -8,6 +8,7 @@ #pragma once #include +#include #include #include @@ -18,7 +19,8 @@ #include #include -namespace cuopt::mps_parser { +namespace cuopt { +namespace MPS_PARSER_EXPORT mps_parser { /** * @brief Main writer class for MPS files @@ -60,4 +62,5 @@ class mps_writer_t { static data_model_view_t create_view(const mps_data_model_t& model); }; // class mps_writer_t -} // namespace cuopt::mps_parser +} // namespace MPS_PARSER_EXPORT mps_parser +} // namespace cuopt diff --git a/cpp/libmps_parser/include/mps_parser/parser.hpp b/cpp/libmps_parser/include/mps_parser/parser.hpp index c5b21dcb13..a9eee6a500 100644 --- a/cpp/libmps_parser/include/mps_parser/parser.hpp +++ b/cpp/libmps_parser/include/mps_parser/parser.hpp @@ -7,12 +7,14 @@ #pragma once +#include #include #include #include -namespace cuopt::mps_parser { +namespace cuopt { +namespace MPS_PARSER_EXPORT mps_parser { /** * @brief Reads the equation from an MPS or QPS file. @@ -55,4 +57,5 @@ template mps_data_model_t parse_mps_from_string(std::string_view mps_contents, bool fixed_mps_format = false); -} // namespace cuopt::mps_parser +} // namespace MPS_PARSER_EXPORT mps_parser +} // namespace cuopt diff --git a/cpp/libmps_parser/include/mps_parser/writer.hpp b/cpp/libmps_parser/include/mps_parser/writer.hpp index 1bcd5bbbff..f0950e5394 100644 --- a/cpp/libmps_parser/include/mps_parser/writer.hpp +++ b/cpp/libmps_parser/include/mps_parser/writer.hpp @@ -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 */ @@ -8,10 +8,12 @@ #pragma once #include +#include // TODO: we might want to eventually rename libmps_parser to libmps_io // (or libcuopt_io if we want to support other hypothetical formats) -namespace cuopt::mps_parser { +namespace cuopt { +namespace MPS_PARSER_EXPORT mps_parser { /** * @brief Writes the problem to an MPS formatted file @@ -25,4 +27,5 @@ namespace cuopt::mps_parser { template void write_mps(const data_model_view_t& problem, const std::string& mps_file_path); -} // namespace cuopt::mps_parser +} // namespace MPS_PARSER_EXPORT mps_parser +} // namespace cuopt diff --git a/cpp/libmps_parser/src/mps_data_model.cpp b/cpp/libmps_parser/src/mps_data_model.cpp index d552a35273..58884f4f96 100644 --- a/cpp/libmps_parser/src/mps_data_model.cpp +++ b/cpp/libmps_parser/src/mps_data_model.cpp @@ -435,9 +435,9 @@ bool mps_data_model_t::has_quadratic_constraints() const noexcept } // NOTE: Explicitly instantiate all types here in order to avoid linker error -template class mps_data_model_t; +template class MPS_PARSER_EXPORT mps_data_model_t; -template class mps_data_model_t; +template class MPS_PARSER_EXPORT mps_data_model_t; // TODO current raft to cusparse wrappers only support int64_t // can be CUSPARSE_INDEX_16U, CUSPARSE_INDEX_32I, CUSPARSE_INDEX_64I diff --git a/cpp/libmps_parser/src/mps_parser.cpp b/cpp/libmps_parser/src/mps_parser.cpp index c58a843ed5..dd50a787df 100644 --- a/cpp/libmps_parser/src/mps_parser.cpp +++ b/cpp/libmps_parser/src/mps_parser.cpp @@ -1692,8 +1692,8 @@ void mps_parser_t::read_bound_and_value(std::string_view line, } // NOTE: Explicitly instantiate all types here in order to avoid linker error -template class mps_parser_t; +template class MPS_PARSER_EXPORT mps_parser_t; -template class mps_parser_t; +template class MPS_PARSER_EXPORT mps_parser_t; } // namespace cuopt::mps_parser diff --git a/cpp/libmps_parser/src/mps_parser.hpp b/cpp/libmps_parser/src/mps_parser.hpp index f2a9ce14e0..94cd1e7e11 100644 --- a/cpp/libmps_parser/src/mps_parser.hpp +++ b/cpp/libmps_parser/src/mps_parser.hpp @@ -7,6 +7,7 @@ #pragma once +#include #include #include @@ -17,7 +18,8 @@ #include #include -namespace cuopt::mps_parser { +namespace cuopt { +namespace MPS_PARSER_EXPORT mps_parser { /** * @brief Different possible types of 'ROWS' @@ -203,4 +205,5 @@ class mps_parser_t { }; // class mps_parser_t -} // namespace cuopt::mps_parser +} // namespace MPS_PARSER_EXPORT mps_parser +} // namespace cuopt diff --git a/cpp/libmps_parser/src/mps_writer.cpp b/cpp/libmps_parser/src/mps_writer.cpp index b112b53476..0d4d57466c 100644 --- a/cpp/libmps_parser/src/mps_writer.cpp +++ b/cpp/libmps_parser/src/mps_writer.cpp @@ -520,7 +520,7 @@ void mps_writer_t::write(const std::string& mps_file_path) mps_file.close(); } -template class mps_writer_t; -template class mps_writer_t; +template class MPS_PARSER_EXPORT mps_writer_t; +template class MPS_PARSER_EXPORT mps_writer_t; } // namespace cuopt::mps_parser diff --git a/cpp/libmps_parser/src/parser.cpp b/cpp/libmps_parser/src/parser.cpp index 681fddf380..7b750ae249 100644 --- a/cpp/libmps_parser/src/parser.cpp +++ b/cpp/libmps_parser/src/parser.cpp @@ -28,12 +28,13 @@ mps_data_model_t parse_mps_from_string(std::string_view mps_contents, return problem; } -template mps_data_model_t parse_mps(const std::string& mps_file, bool fixed_mps_format); -template mps_data_model_t parse_mps(const std::string& mps_file, - bool fixed_mps_format); -template mps_data_model_t parse_mps_from_string(std::string_view mps_contents, - bool fixed_mps_format); -template mps_data_model_t parse_mps_from_string(std::string_view mps_contents, - bool fixed_mps_format); +template MPS_PARSER_EXPORT mps_data_model_t parse_mps(const std::string& mps_file, + bool fixed_mps_format); +template MPS_PARSER_EXPORT mps_data_model_t parse_mps(const std::string& mps_file, + bool fixed_mps_format); +template MPS_PARSER_EXPORT mps_data_model_t parse_mps_from_string( + std::string_view mps_contents, bool fixed_mps_format); +template MPS_PARSER_EXPORT mps_data_model_t parse_mps_from_string( + std::string_view mps_contents, bool fixed_mps_format); } // namespace cuopt::mps_parser diff --git a/cpp/libmps_parser/src/writer.cpp b/cpp/libmps_parser/src/writer.cpp index dab7664f49..25d63ee8d9 100644 --- a/cpp/libmps_parser/src/writer.cpp +++ b/cpp/libmps_parser/src/writer.cpp @@ -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 */ @@ -18,9 +18,9 @@ void write_mps(const data_model_view_t& problem, const std::string& mp writer.write(mps_file_path); } -template void write_mps(const data_model_view_t& problem, - const std::string& mps_file_path); -template void write_mps(const data_model_view_t& problem, - const std::string& mps_file_path); +template MPS_PARSER_EXPORT void write_mps(const data_model_view_t& problem, + const std::string& mps_file_path); +template MPS_PARSER_EXPORT void write_mps( + const data_model_view_t& problem, const std::string& mps_file_path); } // namespace cuopt::mps_parser diff --git a/cpp/src/grpc/client/grpc_client.hpp b/cpp/src/grpc/client/grpc_client.hpp index 4c3aa0c3d3..396e4b3e44 100644 --- a/cpp/src/grpc/client/grpc_client.hpp +++ b/cpp/src/grpc/client/grpc_client.hpp @@ -5,6 +5,7 @@ #pragma once +#include #include #include #include @@ -35,7 +36,8 @@ class ResultResponse; class SubmitJobRequest; } // namespace cuopt::remote -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { // Forward declarations for test helper functions (implemented in grpc_client.cpp) void grpc_test_inject_mock_stub(class grpc_client_t& client, std::shared_ptr stub); @@ -478,4 +480,5 @@ class grpc_client_t { std::string& job_id_out); }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/src/grpc/grpc_problem_mapper.cpp b/cpp/src/grpc/grpc_problem_mapper.cpp index 14461a5a7c..de1fc2a409 100644 --- a/cpp/src/grpc/grpc_problem_mapper.cpp +++ b/cpp/src/grpc/grpc_problem_mapper.cpp @@ -698,10 +698,12 @@ std::vector build_array_chunk_requests( // Explicit template instantiations #if CUOPT_INSTANTIATE_FLOAT -template void map_problem_to_proto(const cpu_optimization_problem_t& cpu_problem, - cuopt::remote::OptimizationProblem* pb_problem); -template void map_proto_to_problem(const cuopt::remote::OptimizationProblem& pb_problem, - cpu_optimization_problem_t& cpu_problem); +template CUOPT_EXPORT void map_problem_to_proto( + const cpu_optimization_problem_t& cpu_problem, + cuopt::remote::OptimizationProblem* pb_problem); +template CUOPT_EXPORT void map_proto_to_problem( + const cuopt::remote::OptimizationProblem& pb_problem, + cpu_optimization_problem_t& cpu_problem); template size_t estimate_problem_proto_size( const cpu_optimization_problem_t& cpu_problem); template void populate_chunked_header_lp( @@ -727,10 +729,12 @@ template std::vector build_array_chunk_req #endif #if CUOPT_INSTANTIATE_DOUBLE -template void map_problem_to_proto(const cpu_optimization_problem_t& cpu_problem, - cuopt::remote::OptimizationProblem* pb_problem); -template void map_proto_to_problem(const cuopt::remote::OptimizationProblem& pb_problem, - cpu_optimization_problem_t& cpu_problem); +template CUOPT_EXPORT void map_problem_to_proto( + const cpu_optimization_problem_t& cpu_problem, + cuopt::remote::OptimizationProblem* pb_problem); +template CUOPT_EXPORT void map_proto_to_problem( + const cuopt::remote::OptimizationProblem& pb_problem, + cpu_optimization_problem_t& cpu_problem); template size_t estimate_problem_proto_size( const cpu_optimization_problem_t& cpu_problem); template void populate_chunked_header_lp( @@ -742,10 +746,10 @@ template void populate_chunked_header_mip( const mip_solver_settings_t& settings, bool enable_incumbents, cuopt::remote::ChunkedProblemHeader* header); -template void map_chunked_header_to_problem( +template CUOPT_EXPORT void map_chunked_header_to_problem( const cuopt::remote::ChunkedProblemHeader& header, cpu_optimization_problem_t& cpu_problem); -template void map_chunked_arrays_to_problem( +template CUOPT_EXPORT void map_chunked_arrays_to_problem( const cuopt::remote::ChunkedProblemHeader& header, const std::map>& arrays, cpu_optimization_problem_t& cpu_problem); diff --git a/cpp/src/grpc/grpc_problem_mapper.hpp b/cpp/src/grpc/grpc_problem_mapper.hpp index db113e2502..aa4396a9f5 100644 --- a/cpp/src/grpc/grpc_problem_mapper.hpp +++ b/cpp/src/grpc/grpc_problem_mapper.hpp @@ -5,6 +5,8 @@ #pragma once +#include + #include #include @@ -18,7 +20,8 @@ namespace cuopt::remote { class ChunkedProblemHeader; } -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { // Forward declarations template @@ -128,4 +131,5 @@ std::vector build_array_chunk_requests( const std::string& upload_id, int64_t chunk_size_bytes); -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/src/grpc/grpc_service_mapper.cpp b/cpp/src/grpc/grpc_service_mapper.cpp index dd969c1587..3b8c79904c 100644 --- a/cpp/src/grpc/grpc_service_mapper.cpp +++ b/cpp/src/grpc/grpc_service_mapper.cpp @@ -69,20 +69,20 @@ cuopt::remote::SubmitJobRequest build_mip_submit_request( // Explicit template instantiations #if CUOPT_INSTANTIATE_FLOAT -template cuopt::remote::SubmitJobRequest build_lp_submit_request( +template CUOPT_EXPORT cuopt::remote::SubmitJobRequest build_lp_submit_request( const cpu_optimization_problem_t& cpu_problem, const pdlp_solver_settings_t& settings); -template cuopt::remote::SubmitJobRequest build_mip_submit_request( +template CUOPT_EXPORT cuopt::remote::SubmitJobRequest build_mip_submit_request( const cpu_optimization_problem_t& cpu_problem, const mip_solver_settings_t& settings, bool enable_incumbents); #endif #if CUOPT_INSTANTIATE_DOUBLE -template cuopt::remote::SubmitJobRequest build_lp_submit_request( +template CUOPT_EXPORT cuopt::remote::SubmitJobRequest build_lp_submit_request( const cpu_optimization_problem_t& cpu_problem, const pdlp_solver_settings_t& settings); -template cuopt::remote::SubmitJobRequest build_mip_submit_request( +template CUOPT_EXPORT cuopt::remote::SubmitJobRequest build_mip_submit_request( const cpu_optimization_problem_t& cpu_problem, const mip_solver_settings_t& settings, bool enable_incumbents); diff --git a/cpp/src/grpc/grpc_service_mapper.hpp b/cpp/src/grpc/grpc_service_mapper.hpp index ed438a0551..8dd1d15c88 100644 --- a/cpp/src/grpc/grpc_service_mapper.hpp +++ b/cpp/src/grpc/grpc_service_mapper.hpp @@ -5,12 +5,15 @@ #pragma once +#include + #include #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { // Forward declarations template @@ -105,4 +108,5 @@ inline cuopt::remote::DeleteRequest build_delete_request(const std::string& job_ return request; } -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/src/grpc/grpc_settings_mapper.cpp b/cpp/src/grpc/grpc_settings_mapper.cpp index 9b503b388e..1369e8fb48 100644 --- a/cpp/src/grpc/grpc_settings_mapper.cpp +++ b/cpp/src/grpc/grpc_settings_mapper.cpp @@ -292,25 +292,33 @@ void map_proto_to_mip_settings(const cuopt::remote::MIPSolverSettings& pb_settin // Explicit template instantiations #if CUOPT_INSTANTIATE_FLOAT -template void map_pdlp_settings_to_proto(const pdlp_solver_settings_t& settings, - cuopt::remote::PDLPSolverSettings* pb_settings); -template void map_proto_to_pdlp_settings(const cuopt::remote::PDLPSolverSettings& pb_settings, - pdlp_solver_settings_t& settings); -template void map_mip_settings_to_proto(const mip_solver_settings_t& settings, - cuopt::remote::MIPSolverSettings* pb_settings); -template void map_proto_to_mip_settings(const cuopt::remote::MIPSolverSettings& pb_settings, - mip_solver_settings_t& settings); +template CUOPT_EXPORT void map_pdlp_settings_to_proto( + const pdlp_solver_settings_t& settings, + cuopt::remote::PDLPSolverSettings* pb_settings); +template CUOPT_EXPORT void map_proto_to_pdlp_settings( + const cuopt::remote::PDLPSolverSettings& pb_settings, + pdlp_solver_settings_t& settings); +template CUOPT_EXPORT void map_mip_settings_to_proto( + const mip_solver_settings_t& settings, + cuopt::remote::MIPSolverSettings* pb_settings); +template CUOPT_EXPORT void map_proto_to_mip_settings( + const cuopt::remote::MIPSolverSettings& pb_settings, + mip_solver_settings_t& settings); #endif #if CUOPT_INSTANTIATE_DOUBLE -template void map_pdlp_settings_to_proto(const pdlp_solver_settings_t& settings, - cuopt::remote::PDLPSolverSettings* pb_settings); -template void map_proto_to_pdlp_settings(const cuopt::remote::PDLPSolverSettings& pb_settings, - pdlp_solver_settings_t& settings); -template void map_mip_settings_to_proto(const mip_solver_settings_t& settings, - cuopt::remote::MIPSolverSettings* pb_settings); -template void map_proto_to_mip_settings(const cuopt::remote::MIPSolverSettings& pb_settings, - mip_solver_settings_t& settings); +template CUOPT_EXPORT void map_pdlp_settings_to_proto( + const pdlp_solver_settings_t& settings, + cuopt::remote::PDLPSolverSettings* pb_settings); +template CUOPT_EXPORT void map_proto_to_pdlp_settings( + const cuopt::remote::PDLPSolverSettings& pb_settings, + pdlp_solver_settings_t& settings); +template CUOPT_EXPORT void map_mip_settings_to_proto( + const mip_solver_settings_t& settings, + cuopt::remote::MIPSolverSettings* pb_settings); +template CUOPT_EXPORT void map_proto_to_mip_settings( + const cuopt::remote::MIPSolverSettings& pb_settings, + mip_solver_settings_t& settings); #endif } // namespace cuopt::linear_programming diff --git a/cpp/src/grpc/grpc_settings_mapper.hpp b/cpp/src/grpc/grpc_settings_mapper.hpp index 6daf0d052b..9e2b46d38d 100644 --- a/cpp/src/grpc/grpc_settings_mapper.hpp +++ b/cpp/src/grpc/grpc_settings_mapper.hpp @@ -5,11 +5,14 @@ #pragma once +#include + #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { // Forward declarations template @@ -58,4 +61,5 @@ template void map_proto_to_mip_settings(const cuopt::remote::MIPSolverSettings& pb_settings, mip_solver_settings_t& settings); -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/src/grpc/grpc_solution_mapper.cpp b/cpp/src/grpc/grpc_solution_mapper.cpp index 3be106cdca..e9aa7206a8 100644 --- a/cpp/src/grpc/grpc_solution_mapper.cpp +++ b/cpp/src/grpc/grpc_solution_mapper.cpp @@ -662,24 +662,24 @@ void build_mip_solution_proto(const cuopt::remote::ChunkedResultHeader& header, // Explicit template instantiations #if CUOPT_INSTANTIATE_FLOAT -template void map_lp_solution_to_proto(const cpu_lp_solution_t& solution, - cuopt::remote::LPSolution* pb_solution); -template cpu_lp_solution_t map_proto_to_lp_solution( +template CUOPT_EXPORT void map_lp_solution_to_proto( + const cpu_lp_solution_t& solution, cuopt::remote::LPSolution* pb_solution); +template CUOPT_EXPORT cpu_lp_solution_t map_proto_to_lp_solution( const cuopt::remote::LPSolution& pb_solution); -template void map_mip_solution_to_proto(const cpu_mip_solution_t& solution, - cuopt::remote::MIPSolution* pb_solution); -template cpu_mip_solution_t map_proto_to_mip_solution( +template CUOPT_EXPORT void map_mip_solution_to_proto( + const cpu_mip_solution_t& solution, cuopt::remote::MIPSolution* pb_solution); +template CUOPT_EXPORT cpu_mip_solution_t map_proto_to_mip_solution( const cuopt::remote::MIPSolution& pb_solution); template size_t estimate_lp_solution_proto_size(const cpu_lp_solution_t& solution); template size_t estimate_mip_solution_proto_size( const cpu_mip_solution_t& solution); -template void populate_chunked_result_header_lp(const cpu_lp_solution_t& solution, - cuopt::remote::ChunkedResultHeader* header); -template void populate_chunked_result_header_mip(const cpu_mip_solution_t& solution, - cuopt::remote::ChunkedResultHeader* header); -template std::map> collect_lp_solution_arrays( +template CUOPT_EXPORT void populate_chunked_result_header_lp( + const cpu_lp_solution_t& solution, cuopt::remote::ChunkedResultHeader* header); +template CUOPT_EXPORT void populate_chunked_result_header_mip( + const cpu_mip_solution_t& solution, cuopt::remote::ChunkedResultHeader* header); +template CUOPT_EXPORT std::map> collect_lp_solution_arrays( const cpu_lp_solution_t& solution); -template std::map> collect_mip_solution_arrays( +template CUOPT_EXPORT std::map> collect_mip_solution_arrays( const cpu_mip_solution_t& solution); template cpu_lp_solution_t chunked_result_to_lp_solution( const cuopt::remote::ChunkedResultHeader& header, @@ -687,35 +687,35 @@ template cpu_lp_solution_t chunked_result_to_lp_solution( template cpu_mip_solution_t chunked_result_to_mip_solution( const cuopt::remote::ChunkedResultHeader& header, const std::map>& arrays); -template void build_lp_solution_proto( +template CUOPT_EXPORT void build_lp_solution_proto( const cuopt::remote::ChunkedResultHeader& header, const std::map>& arrays, cuopt::remote::LPSolution* proto); -template void build_mip_solution_proto( +template CUOPT_EXPORT void build_mip_solution_proto( const cuopt::remote::ChunkedResultHeader& header, const std::map>& arrays, cuopt::remote::MIPSolution* proto); #endif #if CUOPT_INSTANTIATE_DOUBLE -template void map_lp_solution_to_proto(const cpu_lp_solution_t& solution, - cuopt::remote::LPSolution* pb_solution); -template cpu_lp_solution_t map_proto_to_lp_solution( +template CUOPT_EXPORT void map_lp_solution_to_proto( + const cpu_lp_solution_t& solution, cuopt::remote::LPSolution* pb_solution); +template CUOPT_EXPORT cpu_lp_solution_t map_proto_to_lp_solution( const cuopt::remote::LPSolution& pb_solution); -template void map_mip_solution_to_proto(const cpu_mip_solution_t& solution, - cuopt::remote::MIPSolution* pb_solution); -template cpu_mip_solution_t map_proto_to_mip_solution( +template CUOPT_EXPORT void map_mip_solution_to_proto( + const cpu_mip_solution_t& solution, cuopt::remote::MIPSolution* pb_solution); +template CUOPT_EXPORT cpu_mip_solution_t map_proto_to_mip_solution( const cuopt::remote::MIPSolution& pb_solution); template size_t estimate_lp_solution_proto_size(const cpu_lp_solution_t& solution); template size_t estimate_mip_solution_proto_size( const cpu_mip_solution_t& solution); -template void populate_chunked_result_header_lp(const cpu_lp_solution_t& solution, - cuopt::remote::ChunkedResultHeader* header); -template void populate_chunked_result_header_mip( +template CUOPT_EXPORT void populate_chunked_result_header_lp( + const cpu_lp_solution_t& solution, cuopt::remote::ChunkedResultHeader* header); +template CUOPT_EXPORT void populate_chunked_result_header_mip( const cpu_mip_solution_t& solution, cuopt::remote::ChunkedResultHeader* header); -template std::map> collect_lp_solution_arrays( +template CUOPT_EXPORT std::map> collect_lp_solution_arrays( const cpu_lp_solution_t& solution); -template std::map> collect_mip_solution_arrays( +template CUOPT_EXPORT std::map> collect_mip_solution_arrays( const cpu_mip_solution_t& solution); template cpu_lp_solution_t chunked_result_to_lp_solution( const cuopt::remote::ChunkedResultHeader& header, @@ -723,11 +723,11 @@ template cpu_lp_solution_t chunked_result_to_lp_solution( template cpu_mip_solution_t chunked_result_to_mip_solution( const cuopt::remote::ChunkedResultHeader& header, const std::map>& arrays); -template void build_lp_solution_proto( +template CUOPT_EXPORT void build_lp_solution_proto( const cuopt::remote::ChunkedResultHeader& header, const std::map>& arrays, cuopt::remote::LPSolution* proto); -template void build_mip_solution_proto( +template CUOPT_EXPORT void build_mip_solution_proto( const cuopt::remote::ChunkedResultHeader& header, const std::map>& arrays, cuopt::remote::MIPSolution* proto); diff --git a/cpp/src/grpc/grpc_solution_mapper.hpp b/cpp/src/grpc/grpc_solution_mapper.hpp index 127bdb2c96..f73c15bea0 100644 --- a/cpp/src/grpc/grpc_solution_mapper.hpp +++ b/cpp/src/grpc/grpc_solution_mapper.hpp @@ -5,6 +5,8 @@ #pragma once +#include + #include #include @@ -17,7 +19,8 @@ #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { /** * @brief Map cpu_lp_solution_t to protobuf LPSolution message. @@ -180,4 +183,5 @@ void build_mip_solution_proto(const cuopt::remote::ChunkedResultHeader& header, const std::map>& arrays, cuopt::remote::MIPSolution* proto); -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/src/math_optimization/solution_reader.hpp b/cpp/src/math_optimization/solution_reader.hpp index 54eb468f6c..20be0821a5 100644 --- a/cpp/src/math_optimization/solution_reader.hpp +++ b/cpp/src/math_optimization/solution_reader.hpp @@ -1,16 +1,19 @@ /* 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 */ #pragma once +#include + #include #include -namespace cuopt::linear_programming { +namespace cuopt { +namespace CUOPT_EXPORT linear_programming { /** * @brief Reads a solution file and returns the values of specified variables @@ -24,4 +27,5 @@ class solution_reader_t { static std::vector get_variable_values_from_sol_file( const std::string& sol_file_path, const std::vector& variable_names); }; -} // namespace cuopt::linear_programming +} // namespace CUOPT_EXPORT linear_programming +} // namespace cuopt diff --git a/cpp/src/math_optimization/solver_settings.cu b/cpp/src/math_optimization/solver_settings.cu index b968ad18ea..435e41bb46 100644 --- a/cpp/src/math_optimization/solver_settings.cu +++ b/cpp/src/math_optimization/solver_settings.cu @@ -630,25 +630,39 @@ bool solver_settings_t::dump_parameters_to_file(const std::string& pat } #if MIP_INSTANTIATE_FLOAT -template class solver_settings_t; -template void solver_settings_t::set_parameter(const std::string& name, int value); -template void solver_settings_t::set_parameter(const std::string& name, float value); -template void solver_settings_t::set_parameter(const std::string& name, bool value); -template int solver_settings_t::get_parameter(const std::string& name) const; -template float solver_settings_t::get_parameter(const std::string& name) const; -template bool solver_settings_t::get_parameter(const std::string& name) const; -template std::string solver_settings_t::get_parameter(const std::string& name) const; +template class CUOPT_EXPORT solver_settings_t; +template CUOPT_EXPORT void solver_settings_t::set_parameter(const std::string& name, + int value); +template CUOPT_EXPORT void solver_settings_t::set_parameter(const std::string& name, + float value); +template CUOPT_EXPORT void solver_settings_t::set_parameter(const std::string& name, + bool value); +template CUOPT_EXPORT int solver_settings_t::get_parameter( + const std::string& name) const; +template CUOPT_EXPORT float solver_settings_t::get_parameter( + const std::string& name) const; +template CUOPT_EXPORT bool solver_settings_t::get_parameter( + const std::string& name) const; +template CUOPT_EXPORT std::string solver_settings_t::get_parameter( + const std::string& name) const; #endif #if MIP_INSTANTIATE_DOUBLE -template class solver_settings_t; -template void solver_settings_t::set_parameter(const std::string& name, int value); -template void solver_settings_t::set_parameter(const std::string& name, double value); -template void solver_settings_t::set_parameter(const std::string& name, bool value); -template int solver_settings_t::get_parameter(const std::string& name) const; -template double solver_settings_t::get_parameter(const std::string& name) const; -template bool solver_settings_t::get_parameter(const std::string& name) const; -template std::string solver_settings_t::get_parameter(const std::string& name) const; +template class CUOPT_EXPORT solver_settings_t; +template CUOPT_EXPORT void solver_settings_t::set_parameter(const std::string& name, + int value); +template CUOPT_EXPORT void solver_settings_t::set_parameter(const std::string& name, + double value); +template CUOPT_EXPORT void solver_settings_t::set_parameter(const std::string& name, + bool value); +template CUOPT_EXPORT int solver_settings_t::get_parameter( + const std::string& name) const; +template CUOPT_EXPORT double solver_settings_t::get_parameter( + const std::string& name) const; +template CUOPT_EXPORT bool solver_settings_t::get_parameter( + const std::string& name) const; +template CUOPT_EXPORT std::string solver_settings_t::get_parameter( + const std::string& name) const; #endif } // namespace cuopt::linear_programming diff --git a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu index 6b440aed4f..101ef1af75 100644 --- a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu +++ b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu @@ -67,33 +67,39 @@ fj_t::fj_t(mip_solver_context_t& context_, fj_settings_t in_ work_ids_for_related_vars(pb_ptr->n_variables, pb_ptr->handle_ptr->get_stream()) { setval_launch_dims = get_launch_dims_max_occupancy( - (void*)update_assignment_kernel, TPB_setval, pb_ptr->handle_ptr); + update_assignment_kernel_ptr(), TPB_setval, pb_ptr->handle_ptr); update_changed_constraints_launch_dims = - get_launch_dims_max_occupancy((void*)update_changed_constraints_kernel, + get_launch_dims_max_occupancy(update_changed_constraints_kernel_ptr(), TPB_update_changed_constraints, pb_ptr->handle_ptr); resetmoves_launch_dims = get_launch_dims_max_occupancy( - (void*)compute_mtm_moves_kernel, + compute_mtm_moves_kernel_ptr(), TPB_resetmoves, pb_ptr->handle_ptr); resetmoves_bin_launch_dims = get_launch_dims_max_occupancy( - (void*)compute_mtm_moves_kernel, + compute_mtm_moves_kernel_ptr(), TPB_resetmoves, pb_ptr->handle_ptr); update_weights_launch_dims = get_launch_dims_max_occupancy( - (void*)handle_local_minimum_kernel, TPB_localmin, pb_ptr->handle_ptr); + handle_local_minimum_kernel_ptr(), TPB_localmin, pb_ptr->handle_ptr); lift_move_launch_dims = get_launch_dims_max_occupancy( - (void*)update_lift_moves_kernel, TPB_liftmoves, pb_ptr->handle_ptr); - load_balancing_workid_map_launch_dims = get_launch_dims_max_occupancy( - (void*)load_balancing_compute_workid_mappings, TPB_loadbalance, pb_ptr->handle_ptr); - load_balancing_binary_launch_dims = get_launch_dims_max_occupancy( - (void*)load_balancing_compute_scores_binary, TPB_loadbalance, pb_ptr->handle_ptr); - load_balancing_mtm_compute_candidates_launch_dims = get_launch_dims_max_occupancy( - (void*)load_balancing_mtm_compute_candidates, TPB_loadbalance, pb_ptr->handle_ptr); + update_lift_moves_kernel_ptr(), TPB_liftmoves, pb_ptr->handle_ptr); + load_balancing_workid_map_launch_dims = + get_launch_dims_max_occupancy(load_balancing_compute_workid_mappings_kernel_ptr(), + TPB_loadbalance, + pb_ptr->handle_ptr); + load_balancing_binary_launch_dims = + get_launch_dims_max_occupancy(load_balancing_compute_scores_binary_kernel_ptr(), + TPB_loadbalance, + pb_ptr->handle_ptr); + load_balancing_mtm_compute_candidates_launch_dims = + get_launch_dims_max_occupancy(load_balancing_mtm_compute_candidates_kernel_ptr(), + TPB_loadbalance, + pb_ptr->handle_ptr); load_balancing_mtm_compute_scores_launch_dims = get_launch_dims_max_occupancy( - (void*)load_balancing_mtm_compute_scores, TPB_loadbalance, pb_ptr->handle_ptr); + load_balancing_mtm_compute_scores_kernel_ptr(), TPB_loadbalance, pb_ptr->handle_ptr); load_balancing_prepare_launch_dims = get_launch_dims_max_occupancy( - (void*)load_balancing_prepare_iteration, TPB_loadbalance, pb_ptr->handle_ptr); + load_balancing_prepare_iteration_kernel_ptr(), TPB_loadbalance, pb_ptr->handle_ptr); reset_weights(pb_ptr->handle_ptr->get_stream()); // ensure the problem and its transpose are in a valid state (assert checks) @@ -435,7 +441,7 @@ void fj_t::climber_init(i_t climber_idx, const rmm::cuda_stream_view& climber->saved_solution_objective.set_value_async(inf, climber_stream); climber->violation_score.set_value_to_zero_async(climber_stream); climber->weighted_violation_score.set_value_to_zero_async(climber_stream); - init_lhs_and_violation<<<256, 256, 0, climber_stream.value()>>>(view); + launch_init_lhs_and_violation(256, 256, view, climber_stream); // initialize the best_objective values according to the initial assignment f_t best_obj = compute_objective_from_vec( @@ -535,21 +541,32 @@ void fj_t::climber_init(i_t climber_idx, const rmm::cuda_stream_view& // compute the explicit csr_offset to var_idx array if (pb_ptr->binary_indices.size() > 0) - load_balancing_compute_workid_mappings<<<4096, 128, 0, climber_stream.value()>>>( - view, view.row_size_bin_prefix_sum, view.pb.binary_indices, view.work_id_to_bin_var_idx); + launch_load_balancing_compute_workid_mappings(4096, + 128, + view, + view.row_size_bin_prefix_sum, + view.pb.binary_indices, + view.work_id_to_bin_var_idx, + climber_stream); if (pb_ptr->nonbinary_indices.size() > 0) - load_balancing_compute_workid_mappings - <<<4096, 128, 0, climber_stream.value()>>>(view, - view.row_size_nonbin_prefix_sum, - view.pb.nonbinary_indices, - view.work_id_to_nonbin_var_idx); + launch_load_balancing_compute_workid_mappings(4096, + 128, + view, + view.row_size_nonbin_prefix_sum, + view.pb.nonbinary_indices, + view.work_id_to_nonbin_var_idx, + climber_stream); if (pb_ptr->binary_indices.size() > 0) - load_balancing_init_cstr_bounds_csr<<<4096, 128, 0, climber_stream.value()>>>( - view, view.row_size_bin_prefix_sum, view.work_id_to_bin_var_idx); + launch_load_balancing_init_cstr_bounds_csr( + 4096, 128, view, view.row_size_bin_prefix_sum, view.work_id_to_bin_var_idx, climber_stream); if (pb_ptr->nonbinary_indices.size() > 0) - load_balancing_init_cstr_bounds_csr<<<4096, 128, 0, climber_stream.value()>>>( - view, view.row_size_nonbin_prefix_sum, view.work_id_to_nonbin_var_idx); + launch_load_balancing_init_cstr_bounds_csr(4096, + 128, + view, + view.row_size_nonbin_prefix_sum, + view.work_id_to_nonbin_var_idx, + climber_stream); cuopt_assert( pb_ptr->binary_indices.size() + pb_ptr->nonbinary_indices.size() == pb_ptr->n_variables, @@ -599,36 +616,31 @@ void fj_t::load_balancing_score_update(const rmm::cuda_stream_view& st data.iteration_related_variables.clear(stream); - void* kernel_args[] = {&v}; - cudaLaunchCooperativeKernel((void*)load_balancing_prepare_iteration, - grid_load_balancing_prepare, - blocks_load_balancing_prepare, - kernel_args, - 0, - stream); + launch_load_balancing_prepare_iteration( + grid_load_balancing_prepare, blocks_load_balancing_prepare, &v, stream); data.load_balancing_start_event.record(stream); if (pb_ptr->binary_indices.size() > 0) { data.load_balancing_start_event.stream_wait(data.load_balancing_bin_stream.view()); // compute the scores for binary variables (unique delta) - load_balancing_compute_scores_binary<<>>(v); + launch_load_balancing_compute_scores_binary(grid_load_balancing_binary, + blocks_load_balancing_binary, + v, + data.load_balancing_bin_stream.view()); data.load_balancing_bin_finished_event.record(data.load_balancing_bin_stream.view()); } if (pb_ptr->nonbinary_indices.size() > 0) { data.load_balancing_start_event.stream_wait(data.load_balancing_nonbin_stream.view()); - load_balancing_mtm_compute_candidates - <<>>(v); - load_balancing_mtm_compute_scores<<>>(v); + launch_load_balancing_mtm_compute_candidates( + grid_load_balancing_mtm_compute_candidates, + blocks_load_balancing_mtm_compute_candidates, + v, + data.load_balancing_nonbin_stream.view()); + launch_load_balancing_mtm_compute_scores(grid_load_balancing_mtm_compute_scores, + blocks_load_balancing_mtm_compute_scores, + v, + data.load_balancing_nonbin_stream.view()); data.load_balancing_nonbin_finished_event.record(data.load_balancing_nonbin_stream.view()); } @@ -683,11 +695,8 @@ void fj_t::run_step_device(const rmm::cuda_stream_view& climber_stream if (settings.mode == fj_mode_t::ROUNDING) { use_load_balancing = false; } cudaGraph_t graph; - void* kernel_args[] = {&v}; - bool force_reset = false; - void* reset_moves_args[] = {&v, &force_reset}; - bool ignore_load_balancing = false; - void* update_assignment_args[] = {&v, &ignore_load_balancing}; + bool force_reset = false; + bool ignore_load_balancing = false; if (!graph_created || !use_graph) { // CUB temp storage initialization size_t compaction_temp_storage_bytes = 0; @@ -721,52 +730,28 @@ void fj_t::run_step_device(const rmm::cuda_stream_view& climber_stream load_balancing_score_update(climber_stream, climber_idx); } else { if (is_binary_pb) { - RAFT_CUDA_TRY(cudaLaunchCooperativeKernel( - (void*)compute_mtm_moves_kernel, - grid_resetmoves_bin, - blocks_resetmoves_bin, - reset_moves_args, - 0, - climber_stream)); + RAFT_CUDA_TRY( + (launch_compute_mtm_moves_kernel( + grid_resetmoves_bin, blocks_resetmoves_bin, &v, &force_reset, climber_stream))); } else { - RAFT_CUDA_TRY(cudaLaunchCooperativeKernel( - (void*)compute_mtm_moves_kernel, - grid_resetmoves, - blocks_resetmoves, - reset_moves_args, - 0, - climber_stream)); + RAFT_CUDA_TRY( + (launch_compute_mtm_moves_kernel( + grid_resetmoves, blocks_resetmoves, &v, &force_reset, climber_stream))); } } #if FJ_DEBUG_LOAD_BALANCING if (use_load_balancing) { - RAFT_CUDA_TRY(cudaLaunchCooperativeKernel((void*)compute_mtm_moves_kernel, - grid_resetmoves_bin, - blocks_resetmoves_bin, - reset_moves_args, - 0, - climber_stream)); - RAFT_CUDA_TRY(cudaLaunchCooperativeKernel((void*)load_balancing_sanity_checks, - 512, - 128, - kernel_args, - 0, - climber_stream)); + RAFT_CUDA_TRY((launch_compute_mtm_moves_kernel( + grid_resetmoves_bin, blocks_resetmoves_bin, &v, &force_reset, climber_stream))); + RAFT_CUDA_TRY( + (launch_load_balancing_sanity_checks(512, 128, &v, climber_stream))); } #endif - RAFT_CUDA_TRY(cudaLaunchKernel((void*)update_lift_moves_kernel, - grid_lift_move, - blocks_lift_move, - kernel_args, - 0, - climber_stream)); - RAFT_CUDA_TRY(cudaLaunchKernel((void*)update_breakthrough_moves_kernel, - grid_lift_move, - blocks_lift_move, - kernel_args, - 0, - climber_stream)); + RAFT_CUDA_TRY((launch_update_lift_moves_kernel( + grid_lift_move, blocks_lift_move, &v, climber_stream))); + RAFT_CUDA_TRY((launch_update_breakthrough_moves_kernel( + grid_lift_move, blocks_lift_move, &v, climber_stream))); } // compaction kernel @@ -779,32 +764,16 @@ void fj_t::run_step_device(const rmm::cuda_stream_view& climber_stream pb_ptr->n_variables, climber_stream); - RAFT_CUDA_TRY(cudaLaunchKernel((void*)select_variable_kernel, - dim3(1), - dim3(256), - kernel_args, - 0, - climber_stream)); - - RAFT_CUDA_TRY(cudaLaunchCooperativeKernel((void*)handle_local_minimum_kernel, - grid_update_weights, - blocks_update_weights, - kernel_args, - 0, - climber_stream)); + RAFT_CUDA_TRY( + (launch_select_variable_kernel(dim3(1), dim3(256), &v, climber_stream))); + + RAFT_CUDA_TRY((launch_handle_local_minimum_kernel( + grid_update_weights, blocks_update_weights, &v, climber_stream))); raft::copy(data.break_condition.data(), data.temp_break_condition.data(), 1, climber_stream); - RAFT_CUDA_TRY(cudaLaunchKernel((void*)update_assignment_kernel, - grid_setval, - blocks_setval, - update_assignment_args, - 0, - climber_stream)); - RAFT_CUDA_TRY(cudaLaunchKernel((void*)update_changed_constraints_kernel, - 1, - blocks_update_changed_constraints, - kernel_args, - 0, - climber_stream)); + RAFT_CUDA_TRY((launch_update_assignment_kernel( + grid_setval, blocks_setval, &v, &ignore_load_balancing, climber_stream))); + RAFT_CUDA_TRY((launch_update_changed_constraints_kernel( + 1, blocks_update_changed_constraints, &v, climber_stream))); } if (use_graph) { @@ -851,7 +820,7 @@ void fj_t::refresh_lhs_and_violation(const rmm::cuda_stream_view& stre data.violated_constraints.clear(stream); data.violation_score.set_value_to_zero_async(stream); data.weighted_violation_score.set_value_to_zero_async(stream); - init_lhs_and_violation<<<4096, 256, 0, stream>>>(v); + launch_init_lhs_and_violation(4096, 256, v, stream); } template @@ -928,7 +897,7 @@ i_t fj_t::host_loop(solution_t& solution, i_t climber_idx) i_t iterations = data.iterations.value(climber_stream); // make sure we have the current incumbent saved (e.g. in the case of a timeout) - update_best_solution_kernel<<<1, blocks_resetmoves, 0, climber_stream>>>(v); + launch_update_best_solution_kernel(1, blocks_resetmoves, v, climber_stream); // check feasibility with the relative tolerance rather than the violation score raft::copy(solution.assignment.data(), data.best_assignment.data(), 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 e9137503a5..37b113e4b2 100644 --- a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cu +++ b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cu @@ -1438,55 +1438,397 @@ __global__ void handle_local_minimum_kernel(typename fj_t::climber_dat // to save from compilation time, separate those and instantiate separately rather being part of a // class -#define INSTANTIATE(F_TYPE) \ - template __global__ void compute_iteration_related_variables_kernel( \ - typename fj_t::climber_data_t::view_t fj); \ - template __global__ void load_balancing_prepare_iteration( \ - const __grid_constant__ typename fj_t::climber_data_t::view_t fj); \ - template __global__ void load_balancing_compute_workid_mappings( \ - typename fj_t::climber_data_t::view_t fj, \ - raft::device_span row_size_prefix_sum, \ - raft::device_span var_indices, \ - raft::device_span work_id_to_var_idx); \ - template __global__ void load_balancing_init_cstr_bounds_csr( \ - typename fj_t::climber_data_t::view_t fj, \ - raft::device_span row_size_prefix_sum, \ - raft::device_span work_id_to_var_idx); \ - template __global__ void load_balancing_compute_scores_binary( \ - const __grid_constant__ typename fj_t::climber_data_t::view_t fj); \ - template __global__ void load_balancing_mtm_compute_candidates( \ - const __grid_constant__ typename fj_t::climber_data_t::view_t fj); \ - template __global__ void load_balancing_mtm_compute_scores( \ - const __grid_constant__ typename fj_t::climber_data_t::view_t fj); \ - template __global__ void load_balancing_sanity_checks( \ - const __grid_constant__ typename fj_t::climber_data_t::view_t fj); \ - template __global__ void init_lhs_and_violation( \ - typename fj_t::climber_data_t::view_t fj); \ - template __global__ void update_lift_moves_kernel( \ - typename fj_t::climber_data_t::view_t fj); \ - template __global__ void update_breakthrough_moves_kernel( \ - typename fj_t::climber_data_t::view_t fj); \ - template __global__ void handle_local_minimum_kernel( \ - typename fj_t::climber_data_t::view_t fj); \ - template __global__ void update_assignment_kernel( \ - typename fj_t::climber_data_t::view_t fj, bool IgnoreLoadBalancing); \ - template __global__ void update_changed_constraints_kernel( \ - typename fj_t::climber_data_t::view_t fj); \ - template __global__ void update_best_solution_kernel( \ - typename fj_t::climber_data_t::view_t fj); \ - template __global__ void \ - compute_mtm_moves_kernel( \ - typename fj_t::climber_data_t::view_t fj, bool); \ - template __global__ void \ - compute_mtm_moves_kernel( \ - typename fj_t::climber_data_t::view_t fj, bool); \ - template __global__ void \ - compute_mtm_moves_kernel( \ - typename fj_t::climber_data_t::view_t fj, bool); \ - template __global__ void \ - compute_mtm_moves_kernel( \ - typename fj_t::climber_data_t::view_t fj, bool); \ - template __global__ void select_variable_kernel( \ +template +void* update_assignment_kernel_ptr() +{ + return reinterpret_cast(update_assignment_kernel); +} + +template +void* update_changed_constraints_kernel_ptr() +{ + return reinterpret_cast(update_changed_constraints_kernel); +} + +template +void* compute_mtm_moves_kernel_ptr() +{ + return reinterpret_cast(compute_mtm_moves_kernel); +} + +template +void* handle_local_minimum_kernel_ptr() +{ + return reinterpret_cast(handle_local_minimum_kernel); +} + +template +void* update_lift_moves_kernel_ptr() +{ + return reinterpret_cast(update_lift_moves_kernel); +} + +template +void* load_balancing_compute_workid_mappings_kernel_ptr() +{ + return reinterpret_cast(load_balancing_compute_workid_mappings); +} + +template +void* load_balancing_compute_scores_binary_kernel_ptr() +{ + return reinterpret_cast(load_balancing_compute_scores_binary); +} + +template +void* load_balancing_mtm_compute_candidates_kernel_ptr() +{ + return reinterpret_cast(load_balancing_mtm_compute_candidates); +} + +template +void* load_balancing_mtm_compute_scores_kernel_ptr() +{ + return reinterpret_cast(load_balancing_mtm_compute_scores); +} + +template +void* load_balancing_prepare_iteration_kernel_ptr() +{ + return reinterpret_cast(load_balancing_prepare_iteration); +} + +template +cudaError_t launch_load_balancing_prepare_iteration( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream) +{ + void* kernel_args[] = {fj}; + return cudaLaunchCooperativeKernel( + reinterpret_cast(load_balancing_prepare_iteration), + grid, + block, + kernel_args, + 0, + stream); +} + +template +void launch_load_balancing_compute_scores_binary(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + rmm::cuda_stream_view stream) +{ + load_balancing_compute_scores_binary<<>>(fj); +} + +template +void launch_load_balancing_mtm_compute_candidates( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + rmm::cuda_stream_view stream) +{ + load_balancing_mtm_compute_candidates<<>>(fj); +} + +template +void launch_load_balancing_mtm_compute_scores(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + rmm::cuda_stream_view stream) +{ + load_balancing_mtm_compute_scores<<>>(fj); +} + +template +cudaError_t launch_update_lift_moves_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream) +{ + void* kernel_args[] = {fj}; + return cudaLaunchKernel(reinterpret_cast(update_lift_moves_kernel), + grid, + block, + kernel_args, + 0, + stream); +} + +template +cudaError_t launch_update_breakthrough_moves_kernel( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream) +{ + void* kernel_args[] = {fj}; + return cudaLaunchKernel(reinterpret_cast(update_breakthrough_moves_kernel), + grid, + block, + kernel_args, + 0, + stream); +} + +template +cudaError_t launch_compute_mtm_moves_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + bool* force_refresh, + rmm::cuda_stream_view stream) +{ + void* kernel_args[] = {fj, force_refresh}; + return cudaLaunchCooperativeKernel( + reinterpret_cast(compute_mtm_moves_kernel), + grid, + block, + kernel_args, + 0, + stream); +} + +template +cudaError_t launch_select_variable_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream) +{ + void* kernel_args[] = {fj}; + return cudaLaunchKernel( + reinterpret_cast(select_variable_kernel), grid, block, kernel_args, 0, stream); +} + +template +cudaError_t launch_handle_local_minimum_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream) +{ + void* kernel_args[] = {fj}; + return cudaLaunchCooperativeKernel(reinterpret_cast(handle_local_minimum_kernel), + grid, + block, + kernel_args, + 0, + stream); +} + +template +cudaError_t launch_update_assignment_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + bool* ignore_load_balancing, + rmm::cuda_stream_view stream) +{ + void* kernel_args[] = {fj, ignore_load_balancing}; + return cudaLaunchKernel(reinterpret_cast(update_assignment_kernel), + grid, + block, + kernel_args, + 0, + stream); +} + +template +cudaError_t launch_update_changed_constraints_kernel( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream) +{ + void* kernel_args[] = {fj}; + return cudaLaunchKernel(reinterpret_cast(update_changed_constraints_kernel), + grid, + block, + kernel_args, + 0, + stream); +} + +template +cudaError_t launch_load_balancing_sanity_checks(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream) +{ + void* kernel_args[] = {fj}; + return cudaLaunchCooperativeKernel( + reinterpret_cast(load_balancing_sanity_checks), + grid, + block, + kernel_args, + 0, + stream); +} + +template +void launch_init_lhs_and_violation(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + rmm::cuda_stream_view stream) +{ + init_lhs_and_violation<<>>(fj); +} + +template +void launch_load_balancing_compute_workid_mappings( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + raft::device_span row_size_prefix_sum, + raft::device_span var_indices, + raft::device_span work_id_to_var_idx, + rmm::cuda_stream_view stream) +{ + load_balancing_compute_workid_mappings + <<>>(fj, row_size_prefix_sum, var_indices, work_id_to_var_idx); +} + +template +void launch_load_balancing_init_cstr_bounds_csr( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + raft::device_span row_size_prefix_sum, + raft::device_span work_id_to_var_idx, + rmm::cuda_stream_view stream) +{ + load_balancing_init_cstr_bounds_csr + <<>>(fj, row_size_prefix_sum, work_id_to_var_idx); +} + +template +void launch_update_best_solution_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + rmm::cuda_stream_view stream) +{ + update_best_solution_kernel<<>>(fj); +} + +#define INSTANTIATE(F_TYPE) \ + template void* update_assignment_kernel_ptr(); \ + template void* update_changed_constraints_kernel_ptr(); \ + template void* compute_mtm_moves_kernel_ptr(); \ + template void* compute_mtm_moves_kernel_ptr(); \ + template void* handle_local_minimum_kernel_ptr(); \ + template void* update_lift_moves_kernel_ptr(); \ + template void* load_balancing_compute_workid_mappings_kernel_ptr(); \ + template void* load_balancing_compute_scores_binary_kernel_ptr(); \ + template void* load_balancing_mtm_compute_candidates_kernel_ptr(); \ + template void* load_balancing_mtm_compute_scores_kernel_ptr(); \ + template void* load_balancing_prepare_iteration_kernel_ptr(); \ + template cudaError_t launch_load_balancing_prepare_iteration( \ + dim3, dim3, typename fj_t::climber_data_t::view_t*, rmm::cuda_stream_view); \ + template void launch_load_balancing_compute_scores_binary( \ + dim3, dim3, typename fj_t::climber_data_t::view_t, rmm::cuda_stream_view); \ + template void launch_load_balancing_mtm_compute_candidates( \ + dim3, dim3, typename fj_t::climber_data_t::view_t, rmm::cuda_stream_view); \ + template void launch_load_balancing_mtm_compute_scores( \ + dim3, dim3, typename fj_t::climber_data_t::view_t, rmm::cuda_stream_view); \ + template cudaError_t launch_update_lift_moves_kernel( \ + dim3, dim3, typename fj_t::climber_data_t::view_t*, rmm::cuda_stream_view); \ + template cudaError_t launch_update_breakthrough_moves_kernel( \ + dim3, dim3, typename fj_t::climber_data_t::view_t*, rmm::cuda_stream_view); \ + template cudaError_t \ + launch_compute_mtm_moves_kernel( \ + dim3, \ + dim3, \ + typename fj_t::climber_data_t::view_t*, \ + bool*, \ + rmm::cuda_stream_view); \ + template cudaError_t \ + launch_compute_mtm_moves_kernel( \ + dim3, \ + dim3, \ + typename fj_t::climber_data_t::view_t*, \ + bool*, \ + rmm::cuda_stream_view); \ + template cudaError_t launch_select_variable_kernel( \ + dim3, dim3, typename fj_t::climber_data_t::view_t*, rmm::cuda_stream_view); \ + template cudaError_t launch_handle_local_minimum_kernel( \ + dim3, dim3, typename fj_t::climber_data_t::view_t*, rmm::cuda_stream_view); \ + template cudaError_t launch_update_assignment_kernel( \ + dim3, \ + dim3, \ + typename fj_t::climber_data_t::view_t*, \ + bool*, \ + rmm::cuda_stream_view); \ + template cudaError_t launch_update_changed_constraints_kernel( \ + dim3, dim3, typename fj_t::climber_data_t::view_t*, rmm::cuda_stream_view); \ + template cudaError_t launch_load_balancing_sanity_checks( \ + dim3, dim3, typename fj_t::climber_data_t::view_t*, rmm::cuda_stream_view); \ + template void launch_init_lhs_and_violation( \ + dim3, dim3, typename fj_t::climber_data_t::view_t, rmm::cuda_stream_view); \ + template void launch_load_balancing_compute_workid_mappings( \ + dim3, \ + dim3, \ + typename fj_t::climber_data_t::view_t, \ + raft::device_span, \ + raft::device_span, \ + raft::device_span, \ + rmm::cuda_stream_view); \ + template void launch_load_balancing_init_cstr_bounds_csr( \ + dim3, \ + dim3, \ + typename fj_t::climber_data_t::view_t, \ + raft::device_span, \ + raft::device_span, \ + rmm::cuda_stream_view); \ + template void launch_update_best_solution_kernel( \ + dim3, dim3, typename fj_t::climber_data_t::view_t, rmm::cuda_stream_view); \ + template __global__ void compute_iteration_related_variables_kernel( \ + typename fj_t::climber_data_t::view_t fj); \ + template __global__ void load_balancing_prepare_iteration( \ + const __grid_constant__ typename fj_t::climber_data_t::view_t fj); \ + template __global__ void load_balancing_compute_workid_mappings( \ + typename fj_t::climber_data_t::view_t fj, \ + raft::device_span row_size_prefix_sum, \ + raft::device_span var_indices, \ + raft::device_span work_id_to_var_idx); \ + template __global__ void load_balancing_init_cstr_bounds_csr( \ + typename fj_t::climber_data_t::view_t fj, \ + raft::device_span row_size_prefix_sum, \ + raft::device_span work_id_to_var_idx); \ + template __global__ void load_balancing_compute_scores_binary( \ + const __grid_constant__ typename fj_t::climber_data_t::view_t fj); \ + template __global__ void load_balancing_mtm_compute_candidates( \ + const __grid_constant__ typename fj_t::climber_data_t::view_t fj); \ + template __global__ void load_balancing_mtm_compute_scores( \ + const __grid_constant__ typename fj_t::climber_data_t::view_t fj); \ + template __global__ void load_balancing_sanity_checks( \ + const __grid_constant__ typename fj_t::climber_data_t::view_t fj); \ + template __global__ void init_lhs_and_violation( \ + typename fj_t::climber_data_t::view_t fj); \ + template __global__ void update_lift_moves_kernel( \ + typename fj_t::climber_data_t::view_t fj); \ + template __global__ void update_breakthrough_moves_kernel( \ + typename fj_t::climber_data_t::view_t fj); \ + template __global__ void handle_local_minimum_kernel( \ + typename fj_t::climber_data_t::view_t fj); \ + template __global__ void update_assignment_kernel( \ + typename fj_t::climber_data_t::view_t fj, bool IgnoreLoadBalancing); \ + template __global__ void update_changed_constraints_kernel( \ + typename fj_t::climber_data_t::view_t fj); \ + template __global__ void update_best_solution_kernel( \ + typename fj_t::climber_data_t::view_t fj); \ + template __global__ void \ + compute_mtm_moves_kernel( \ + typename fj_t::climber_data_t::view_t fj, bool); \ + template __global__ void \ + compute_mtm_moves_kernel( \ + typename fj_t::climber_data_t::view_t fj, bool); \ + template __global__ void \ + compute_mtm_moves_kernel( \ + typename fj_t::climber_data_t::view_t fj, bool); \ + template __global__ void \ + compute_mtm_moves_kernel( \ + typename fj_t::climber_data_t::view_t fj, bool); \ + template __global__ void select_variable_kernel( \ typename fj_t::climber_data_t::view_t fj); #if MIP_INSTANTIATE_FLOAT diff --git a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cuh b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cuh index 55fd4e61f1..4ba7d991bf 100644 --- a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cuh +++ b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cuh @@ -21,6 +21,152 @@ enum class weight_strategy_t { Increment, Multiply }; template __global__ void compute_iteration_related_variables_kernel( typename fj_t::climber_data_t::view_t fj); + +template +void* update_assignment_kernel_ptr(); + +template +void* update_changed_constraints_kernel_ptr(); + +template +void* compute_mtm_moves_kernel_ptr(); + +template +void* handle_local_minimum_kernel_ptr(); + +template +void* update_lift_moves_kernel_ptr(); + +template +void* load_balancing_compute_workid_mappings_kernel_ptr(); + +template +void* load_balancing_compute_scores_binary_kernel_ptr(); + +template +void* load_balancing_mtm_compute_candidates_kernel_ptr(); + +template +void* load_balancing_mtm_compute_scores_kernel_ptr(); + +template +void* load_balancing_prepare_iteration_kernel_ptr(); + +template +cudaError_t launch_load_balancing_prepare_iteration( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream); + +template +void launch_load_balancing_compute_scores_binary(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + rmm::cuda_stream_view stream); + +template +void launch_load_balancing_mtm_compute_candidates( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + rmm::cuda_stream_view stream); + +template +void launch_load_balancing_mtm_compute_scores(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + rmm::cuda_stream_view stream); + +template +cudaError_t launch_update_lift_moves_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream); + +template +cudaError_t launch_update_breakthrough_moves_kernel( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream); + +template +cudaError_t launch_compute_mtm_moves_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + bool* force_refresh, + rmm::cuda_stream_view stream); + +template +cudaError_t launch_select_variable_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream); + +template +cudaError_t launch_handle_local_minimum_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream); + +template +cudaError_t launch_update_assignment_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + bool* ignore_load_balancing, + rmm::cuda_stream_view stream); + +template +cudaError_t launch_update_changed_constraints_kernel( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream); + +template +cudaError_t launch_load_balancing_sanity_checks(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t* fj, + rmm::cuda_stream_view stream); + +template +void launch_init_lhs_and_violation(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + rmm::cuda_stream_view stream); + +template +void launch_load_balancing_compute_workid_mappings( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + raft::device_span row_size_prefix_sum, + raft::device_span var_indices, + raft::device_span work_id_to_var_idx, + rmm::cuda_stream_view stream); + +template +void launch_load_balancing_init_cstr_bounds_csr( + dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + raft::device_span row_size_prefix_sum, + raft::device_span work_id_to_var_idx, + rmm::cuda_stream_view stream); + +template +void launch_update_best_solution_kernel(dim3 grid, + dim3 block, + typename fj_t::climber_data_t::view_t fj, + rmm::cuda_stream_view stream); + template __global__ void load_balancing_prepare_iteration( const __grid_constant__ typename fj_t::climber_data_t::view_t fj); diff --git a/cpp/src/mip_heuristics/solve.cu b/cpp/src/mip_heuristics/solve.cu index 2c2c05a795..8c8c007e9e 100644 --- a/cpp/src/mip_heuristics/solve.cu +++ b/cpp/src/mip_heuristics/solve.cu @@ -834,19 +834,19 @@ std::unique_ptr> solve_mip( } #define INSTANTIATE(F_TYPE) \ - template mip_solution_t solve_mip( \ + template CUOPT_EXPORT mip_solution_t solve_mip( \ optimization_problem_t& op_problem, \ mip_solver_settings_t const& settings); \ \ - template mip_solution_t solve_mip( \ + template CUOPT_EXPORT mip_solution_t solve_mip( \ raft::handle_t const* handle_ptr, \ const cuopt::mps_parser::mps_data_model_t& mps_data_model, \ mip_solver_settings_t const& settings); \ \ - template std::unique_ptr> solve_mip( \ + template CUOPT_EXPORT std::unique_ptr> solve_mip( \ cpu_optimization_problem_t&, mip_solver_settings_t const&); \ \ - template std::unique_ptr> solve_mip( \ + template CUOPT_EXPORT std::unique_ptr> solve_mip( \ optimization_problem_interface_t*, mip_solver_settings_t const&); #if MIP_INSTANTIATE_FLOAT diff --git a/cpp/src/mip_heuristics/solver_settings.cu b/cpp/src/mip_heuristics/solver_settings.cu index dededa2f91..1a08d885e8 100644 --- a/cpp/src/mip_heuristics/solver_settings.cu +++ b/cpp/src/mip_heuristics/solver_settings.cu @@ -48,11 +48,11 @@ mip_solver_settings_t::get_tolerances() const noexcept // Explicit template instantiations for common types #if MIP_INSTANTIATE_FLOAT -template class mip_solver_settings_t; +template class CUOPT_EXPORT mip_solver_settings_t; #endif #if MIP_INSTANTIATE_DOUBLE -template class mip_solver_settings_t; +template class CUOPT_EXPORT mip_solver_settings_t; #endif } // namespace cuopt::linear_programming diff --git a/cpp/src/mip_heuristics/solver_solution.cu b/cpp/src/mip_heuristics/solver_solution.cu index 8f6f8de05f..41740d82e5 100644 --- a/cpp/src/mip_heuristics/solver_solution.cu +++ b/cpp/src/mip_heuristics/solver_solution.cu @@ -256,10 +256,10 @@ void mip_solution_t::log_detailed_summary() const } #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT -template class mip_solution_t; +template class CUOPT_EXPORT mip_solution_t; #endif #if MIP_INSTANTIATE_DOUBLE -template class mip_solution_t; +template class CUOPT_EXPORT mip_solution_t; #endif } // namespace cuopt::linear_programming diff --git a/cpp/src/pdlp/cpu_optimization_problem.cpp b/cpp/src/pdlp/cpu_optimization_problem.cpp index de1f74ed47..c1ffe387a1 100644 --- a/cpp/src/pdlp/cpu_optimization_problem.cpp +++ b/cpp/src/pdlp/cpu_optimization_problem.cpp @@ -1058,10 +1058,10 @@ void cpu_optimization_problem_t::copy_variable_types_to_host(var_t* ou // ============================================================================== #if MIP_INSTANTIATE_FLOAT -template class cpu_optimization_problem_t; +template class CUOPT_EXPORT cpu_optimization_problem_t; #endif #if MIP_INSTANTIATE_DOUBLE -template class cpu_optimization_problem_t; +template class CUOPT_EXPORT cpu_optimization_problem_t; #endif } // namespace cuopt::linear_programming diff --git a/cpp/src/pdlp/cpu_pdlp_warm_start_data.cu b/cpp/src/pdlp/cpu_pdlp_warm_start_data.cu index b078bc4779..677682bd48 100644 --- a/cpp/src/pdlp/cpu_pdlp_warm_start_data.cu +++ b/cpp/src/pdlp/cpu_pdlp_warm_start_data.cu @@ -109,17 +109,17 @@ pdlp_warm_start_data_t convert_to_gpu_warmstart( } #if MIP_INSTANTIATE_DOUBLE -template cpu_pdlp_warm_start_data_t convert_to_cpu_warmstart( +template CUOPT_EXPORT cpu_pdlp_warm_start_data_t convert_to_cpu_warmstart( const pdlp_warm_start_data_t&, rmm::cuda_stream_view); -template pdlp_warm_start_data_t convert_to_gpu_warmstart( +template CUOPT_EXPORT pdlp_warm_start_data_t convert_to_gpu_warmstart( const cpu_pdlp_warm_start_data_t&, rmm::cuda_stream_view); #endif #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT -template cpu_pdlp_warm_start_data_t convert_to_cpu_warmstart( +template CUOPT_EXPORT cpu_pdlp_warm_start_data_t convert_to_cpu_warmstart( const pdlp_warm_start_data_t&, rmm::cuda_stream_view); -template pdlp_warm_start_data_t convert_to_gpu_warmstart( +template CUOPT_EXPORT pdlp_warm_start_data_t convert_to_gpu_warmstart( const cpu_pdlp_warm_start_data_t&, rmm::cuda_stream_view); #endif diff --git a/cpp/src/pdlp/optimization_problem.cu b/cpp/src/pdlp/optimization_problem.cu index a6f0d30ea8..64d302a146 100644 --- a/cpp/src/pdlp/optimization_problem.cu +++ b/cpp/src/pdlp/optimization_problem.cu @@ -1593,10 +1593,10 @@ optimization_problem_t optimization_problem_t::convert // ============================================================================== // Explicit template instantiations matching MIP constants #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT -template class optimization_problem_t; +template class CUOPT_EXPORT optimization_problem_t; #endif #if MIP_INSTANTIATE_DOUBLE -template class optimization_problem_t; +template class CUOPT_EXPORT optimization_problem_t; #endif #if PDLP_INSTANTIATE_FLOAT || MIP_INSTANTIATE_FLOAT diff --git a/cpp/src/pdlp/pdlp_warm_start_data.cu b/cpp/src/pdlp/pdlp_warm_start_data.cu index 80abf015d8..1c88a8a0df 100644 --- a/cpp/src/pdlp/pdlp_warm_start_data.cu +++ b/cpp/src/pdlp/pdlp_warm_start_data.cu @@ -179,10 +179,10 @@ void pdlp_warm_start_data_t::check_sizes() } #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT -template class pdlp_warm_start_data_t; +template class CUOPT_EXPORT pdlp_warm_start_data_t; #endif #if MIP_INSTANTIATE_DOUBLE -template class pdlp_warm_start_data_t; +template class CUOPT_EXPORT pdlp_warm_start_data_t; #endif } // namespace cuopt::linear_programming diff --git a/cpp/src/pdlp/solution_conversion.cu b/cpp/src/pdlp/solution_conversion.cu index 8ec3c20b27..b6a57f7947 100644 --- a/cpp/src/pdlp/solution_conversion.cu +++ b/cpp/src/pdlp/solution_conversion.cu @@ -217,9 +217,9 @@ cuopt::cython::mip_ret_t cpu_mip_solution_t::to_cpu_mip_ret_t() // Explicit template instantiations template cuopt::cython::linear_programming_ret_t gpu_lp_solution_t::to_linear_programming_ret_t(); -template cuopt::cython::mip_ret_t gpu_mip_solution_t::to_mip_ret_t(); -template cuopt::cython::linear_programming_ret_t +template CUOPT_EXPORT cuopt::cython::mip_ret_t gpu_mip_solution_t::to_mip_ret_t(); +template CUOPT_EXPORT cuopt::cython::linear_programming_ret_t cpu_lp_solution_t::to_cpu_linear_programming_ret_t(); -template cuopt::cython::mip_ret_t cpu_mip_solution_t::to_cpu_mip_ret_t(); +template CUOPT_EXPORT cuopt::cython::mip_ret_t cpu_mip_solution_t::to_cpu_mip_ret_t(); } // namespace cuopt::linear_programming diff --git a/cpp/src/pdlp/solve.cu b/cpp/src/pdlp/solve.cu index bb2d193e18..4d621b2c7a 100644 --- a/cpp/src/pdlp/solve.cu +++ b/cpp/src/pdlp/solve.cu @@ -2099,28 +2099,28 @@ std::unique_ptr> solve_lp( } #define INSTANTIATE(F_TYPE) \ - template optimization_problem_solution_t solve_lp( \ + template CUOPT_EXPORT optimization_problem_solution_t solve_lp( \ optimization_problem_t& op_problem, \ pdlp_solver_settings_t const& settings, \ bool problem_checking, \ bool use_pdlp_solver_mode, \ bool is_batch_mode); \ \ - template optimization_problem_solution_t solve_lp( \ + template CUOPT_EXPORT optimization_problem_solution_t solve_lp( \ raft::handle_t const* handle_ptr, \ const cuopt::mps_parser::mps_data_model_t& mps_data_model, \ pdlp_solver_settings_t const& settings, \ bool problem_checking, \ bool use_pdlp_solver_mode); \ \ - template std::unique_ptr> solve_lp( \ + template CUOPT_EXPORT std::unique_ptr> solve_lp( \ cpu_optimization_problem_t&, \ pdlp_solver_settings_t const&, \ bool, \ bool, \ bool); \ \ - template std::unique_ptr> solve_lp( \ + template CUOPT_EXPORT std::unique_ptr> solve_lp( \ optimization_problem_interface_t*, \ pdlp_solver_settings_t const&, \ bool, \ diff --git a/cpp/src/pdlp/solver_settings.cu b/cpp/src/pdlp/solver_settings.cu index 28e7428fac..464eb2fba2 100644 --- a/cpp/src/pdlp/solver_settings.cu +++ b/cpp/src/pdlp/solver_settings.cu @@ -415,11 +415,11 @@ pdlp_solver_settings_t::get_pdlp_warm_start_data_view() const noexcept } #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT -template class pdlp_solver_settings_t; +template class CUOPT_EXPORT pdlp_solver_settings_t; #endif #if MIP_INSTANTIATE_DOUBLE -template class pdlp_solver_settings_t; +template class CUOPT_EXPORT pdlp_solver_settings_t; #endif } // namespace cuopt::linear_programming diff --git a/cpp/src/pdlp/solver_solution.cu b/cpp/src/pdlp/solver_solution.cu index ec0492dac3..4af1342e19 100644 --- a/cpp/src/pdlp/solver_solution.cu +++ b/cpp/src/pdlp/solver_solution.cu @@ -453,10 +453,10 @@ void optimization_problem_solution_t::write_to_sol_file( } #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT -template class optimization_problem_solution_t; +template class CUOPT_EXPORT optimization_problem_solution_t; #endif #if MIP_INSTANTIATE_DOUBLE -template class optimization_problem_solution_t; +template class CUOPT_EXPORT optimization_problem_solution_t; #endif } // namespace cuopt::linear_programming diff --git a/cpp/src/routing/ges/compute_fragment_ejections.cu b/cpp/src/routing/ges/compute_fragment_ejections.cu index 89d54eb1a1..f557ed1238 100644 --- a/cpp/src/routing/ges/compute_fragment_ejections.cu +++ b/cpp/src/routing/ges/compute_fragment_ejections.cu @@ -113,6 +113,49 @@ __global__ void kernel_get_best_insertion_ejection_solution( fragment_step); } +template +void* kernel_get_best_insertion_ejection_solution_ptr() +{ + return reinterpret_cast( + kernel_get_best_insertion_ejection_solution); +} + +template +void launch_kernel_get_best_insertion_ejection_solution( + dim3 grid, + dim3 block, + size_t shared_memory, + rmm::cuda_stream_view stream, + typename solution_t::view_t solution, + const request_info_t* request_id, + i_t* p_scores, + i_t fragment_size, + i_t fragment_step, + feasible_move_t feasible_candidates, + int64_t seed) +{ + kernel_get_best_insertion_ejection_solution + <<>>( + solution, request_id, p_scores, fragment_size, fragment_step, feasible_candidates, seed); +} + +#define INSTANTIATE_KERNEL_GET_BEST_INSERTION_EJECTION_SOLUTION(BLOCK_SIZE, REQUEST) \ + template void* \ + kernel_get_best_insertion_ejection_solution_ptr(); \ + template void \ + launch_kernel_get_best_insertion_ejection_solution( \ + dim3, \ + dim3, \ + size_t, \ + rmm::cuda_stream_view, \ + typename solution_t::view_t, \ + const request_info_t*, \ + int*, \ + int, \ + int, \ + feasible_move_t, \ + int64_t); + template __global__ void kernel_get_best_insertion_ejection_solution<32, int, float, request_t::PDP>( typename solution_t::view_t solution, @@ -186,6 +229,17 @@ kernel_get_best_insertion_ejection_solution<512, int, float, request_t::VRP>( int fragment_step, feasible_move_t feasible_candidates, int64_t seed); + +INSTANTIATE_KERNEL_GET_BEST_INSERTION_EJECTION_SOLUTION(32, PDP) +INSTANTIATE_KERNEL_GET_BEST_INSERTION_EJECTION_SOLUTION(64, PDP) +INSTANTIATE_KERNEL_GET_BEST_INSERTION_EJECTION_SOLUTION(128, PDP) +INSTANTIATE_KERNEL_GET_BEST_INSERTION_EJECTION_SOLUTION(512, PDP) +INSTANTIATE_KERNEL_GET_BEST_INSERTION_EJECTION_SOLUTION(32, VRP) +INSTANTIATE_KERNEL_GET_BEST_INSERTION_EJECTION_SOLUTION(64, VRP) +INSTANTIATE_KERNEL_GET_BEST_INSERTION_EJECTION_SOLUTION(128, VRP) +INSTANTIATE_KERNEL_GET_BEST_INSERTION_EJECTION_SOLUTION(512, VRP) + +#undef INSTANTIATE_KERNEL_GET_BEST_INSERTION_EJECTION_SOLUTION } // namespace detail } // namespace routing } // namespace cuopt diff --git a/cpp/src/routing/ges/compute_fragment_ejections.cuh b/cpp/src/routing/ges/compute_fragment_ejections.cuh index 1a53d45674..ff2e851ff4 100644 --- a/cpp/src/routing/ges/compute_fragment_ejections.cuh +++ b/cpp/src/routing/ges/compute_fragment_ejections.cuh @@ -35,6 +35,23 @@ __global__ void kernel_get_best_insertion_ejection_solution( feasible_move_t feasible_candidates, int64_t seed); +template +void* kernel_get_best_insertion_ejection_solution_ptr(); + +template +void launch_kernel_get_best_insertion_ejection_solution( + dim3 grid, + dim3 block, + size_t shared_memory, + rmm::cuda_stream_view stream, + typename solution_t::view_t solution, + const request_info_t* request_id, + i_t* p_scores, + i_t fragment_size, + i_t fragment_step, + feasible_move_t feasible_candidates, + int64_t seed); + template ::execute_best_insertion_ejectio raft::alignTo(fragment_size * request_info_t::size() * sizeof(i_t), sizeof(infeasible_cost_t)); if (!set_shmem_of_kernel( - kernel_get_best_insertion_ejection_solution, + kernel_get_best_insertion_ejection_solution_ptr(), shared_for_delete_array + shared_for_tmp_route)) { return false; } - kernel_get_best_insertion_ejection_solution - <<get_num_requests() * fragment_step, - threads_per_block, - shared_for_delete_array + shared_for_tmp_route, - solution_ptr->sol_handle->get_stream()>>>( - solution_ptr->view(), - d_request, - p_scores_.data(), - fragment_size, - fragment_step, - feasible_move_t(cuopt::make_span(feasible_candidates_data_), - feasible_candidates_size_.data(), - solution_ptr->get_num_orders(), - solution_ptr->problem_ptr->get_max_break_dimensions(), - solution_ptr->get_n_routes()), - seed_generator::get_seed()); + launch_kernel_get_best_insertion_ejection_solution( + solution_ptr->get_num_requests() * fragment_step, + threads_per_block, + shared_for_delete_array + shared_for_tmp_route, + solution_ptr->sol_handle->get_stream(), + solution_ptr->view(), + d_request, + p_scores_.data(), + fragment_size, + fragment_step, + feasible_move_t(cuopt::make_span(feasible_candidates_data_), + feasible_candidates_size_.data(), + solution_ptr->get_num_orders(), + solution_ptr->problem_ptr->get_max_break_dimensions(), + solution_ptr->get_n_routes()), + seed_generator::get_seed()); RAFT_CHECK_CUDA(solution_ptr->sol_handle->get_stream()); } diff --git a/cpp/src/utilities/logger.hpp b/cpp/src/utilities/logger.hpp index 08556a4c78..e14b511c65 100644 --- a/cpp/src/utilities/logger.hpp +++ b/cpp/src/utilities/logger.hpp @@ -1,12 +1,13 @@ /* 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 */ #pragma once +#include #include #include @@ -17,7 +18,7 @@ #include #include -namespace cuopt { +namespace CUOPT_EXPORT cuopt { /** * @brief Get the default logger. @@ -42,4 +43,4 @@ class init_logger_t { init_logger_t(std::string log_file, bool log_to_console); }; -} // namespace cuopt +} // namespace CUOPT_EXPORT cuopt