Skip to content

fix(visibility): add symbol visibility controls to cuvs #2101

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
rapids-bot merged 7 commits into from
May 21, 2026
Merged

fix(visibility): add symbol visibility controls to cuvs #2101

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
7f3c596
fix(visibility): add CUDA_VISIBILITY_PRESET hidden to cuvs_objs
vyasr May 16, 2026
7118e80
fix(visibility): add CUVS_EXPORT to symbols hidden by CUDA_VISIBILITY…
vyasr May 17, 2026
9a6cde7
fix(visibility): fix CUVS_EXPORT placement to be NVCC-compatible
vyasr May 17, 2026
ae178ef
fix(visibility): fix remaining 4 hidden symbol linker errors
vyasr May 17, 2026
c956741
fix(visibility): fix brute_force, cagra, resize_list symbol export
vyasr May 17, 2026
61b1049
fix(visibility): use struct CUVS_EXPORT on definitions instead of ins…
vyasr May 17, 2026
528edaf
fix(visibility): add cuda_fp16.hpp wrapper to fix __half type visibility
vyasr May 19, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions cpp/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1072,6 +1072,7 @@ if(NOT BUILD_CPU_ONLY)
CUDA_STANDARD_REQUIRED ON
POSITION_INDEPENDENT_CODE ON
CXX_VISIBILITY_PRESET hidden
CUDA_VISIBILITY_PRESET hidden
VISIBILITY_INLINES_HIDDEN ON
)
target_compile_options(
Expand Down
18 changes: 18 additions & 0 deletions cpp/include/cuvs/core/cuda_fp16.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,18 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
* SPDX-License-Identifier: Apache-2.0
*/

#pragma once

// Wrapper for cuda_fp16.h that ensures __half gets default symbol visibility.
//
// GCC's "type visibility" rule causes template instantiations over __half to
// inherit hidden visibility when -fvisibility=hidden is in effect, because
// __half is a user-defined type first seen under hidden visibility. By
// including cuda_fp16.h under #pragma GCC visibility push(default), the __half
// type acquires default visibility, and downstream template instantiations
// (e.g., index<__half, ...>) will be properly exported from shared libraries.
#pragma GCC visibility push(default)
#include <cuda_fp16.h> // NOLINT
#pragma GCC visibility pop
2 changes: 1 addition & 1 deletion cpp/include/cuvs/distance/distance.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@
#pragma once

#include <cstdint>
#include <cuda_fp16.h>
#include <cuvs/core/cuda_fp16.hpp>
#include <cuvs/core/export.hpp>
#include <raft/core/device_csr_matrix.hpp>
#include <raft/core/device_mdspan.hpp>
Expand Down
5 changes: 3 additions & 2 deletions cpp/include/cuvs/neighbors/brute_force.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@
#include <raft/core/handle.hpp>
#include <raft/core/host_mdspan.hpp>

#include <cuda_fp16.h>
#include <cuvs/core/cuda_fp16.hpp>
#include <cuvs/core/export.hpp>

namespace CUVS_EXPORT cuvs {
Expand All @@ -35,7 +35,7 @@ struct search_params : cuvs::neighbors::search_params {};
* @tparam T data element type
*/
template <typename T, typename DistT = T>
struct index : cuvs::neighbors::index {
struct CUVS_EXPORT index : cuvs::neighbors::index {
using index_params_type = brute_force::index_params;
using search_params_type = brute_force::search_params;
using index_type = int64_t;
Expand Down Expand Up @@ -166,6 +166,7 @@ struct index : cuvs::neighbors::index {
raft::device_matrix_view<const T, int64_t, raft::row_major> dataset_view_;
DistT metric_arg_;
};

/**
* @}
*/
Expand Down
3 changes: 2 additions & 1 deletion cpp/include/cuvs/neighbors/cagra.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -391,7 +391,7 @@ static_assert(std::is_aggregate_v<search_params>);
*
*/
template <typename T, typename IdxT>
struct index : cuvs::neighbors::index {
struct CUVS_EXPORT index : cuvs::neighbors::index {
using index_params_type = cagra::index_params;
using search_params_type = cagra::search_params;
using index_type = IdxT;
Expand Down Expand Up @@ -883,6 +883,7 @@ struct index : cuvs::neighbors::index {
size_t dim_ = 0;
size_t graph_degree_ = 0;
};

/**
* @}
*/
Expand Down
12 changes: 6 additions & 6 deletions cpp/include/cuvs/neighbors/common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,7 @@

#include <cuvs/core/bitmap.hpp>
#include <cuvs/core/bitset.hpp>
#include <cuvs/core/export.hpp>
#include <raft/core/detail/macros.hpp>

#include <memory>
Expand All @@ -27,7 +28,6 @@

#ifdef __cpp_lib_bitops
#include <bit>
#include <cuvs/core/export.hpp>
#endif

namespace CUVS_EXPORT cuvs {
Expand Down Expand Up @@ -811,11 +811,11 @@ using enable_if_valid_list_t = typename enable_if_valid_list<ListT, T>::type;
* `cuvs::neighbors::ivf_pq::helpers::resize_list` which handle type casting internally.
*/
template <typename ListT>
void resize_list(raft::resources const& res,
std::shared_ptr<ListT>& orig_list, // NOLINT
const typename ListT::spec_type& spec,
typename ListT::size_type new_used_size,
typename ListT::size_type old_used_size);
CUVS_EXPORT void resize_list(raft::resources const& res,

@divyegala divyegala May 21, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do we have to manually mark templated functions for export?

@vyasr vyasr May 21, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think this was basically an nvcc bug/quirk where it wouldn't accept the attribute in the same places that gcc would. Specifically I believe it was when there was a templated function that we wanted to export, the markup on the namespace wasn't affecting the symbol so we had to put it on individual symbols. I'd say more quirk than bug since I don't believe there's any standard requiring the attribute on the namespace to work, gcc just happens to support it.

std::shared_ptr<ListT>& orig_list, // NOLINT
const typename ListT::spec_type& spec,
typename ListT::size_type new_used_size,
typename ListT::size_type old_used_size);

/**
* Serialize a list to an output stream.
Expand Down
2 changes: 1 addition & 1 deletion cpp/include/cuvs/neighbors/composite/index.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,7 +41,7 @@ namespace composite {
* @endcode
*/
template <typename T, typename IdxT, typename OutputIdxT = IdxT>
class composite_index {
class CUVS_EXPORT composite_index {
public:
using value_type = T;
using index_type = IdxT;
Expand Down
1 change: 1 addition & 0 deletions cpp/include/cuvs/neighbors/ivf_flat.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3576,4 +3576,5 @@ __device__ __forceinline__ void compute_dist_udf_impl(AccT& acc, AccT x, AccT y)

} // namespace ivf_flat
} // namespace neighbors

} // namespace CUVS_EXPORT cuvs
2 changes: 1 addition & 1 deletion cpp/include/cuvs/neighbors/ivf_pq.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@

#pragma once

#include <cuda_fp16.h>
#include <cuvs/core/cuda_fp16.hpp>

#include <cuvs/neighbors/common.hpp>

Expand Down
2 changes: 1 addition & 1 deletion cpp/include/cuvs/neighbors/nn_descent.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@

#include <cuvs/distance/distance.hpp>

#include <cuda_fp16.h>
#include <cuvs/core/cuda_fp16.hpp>
#include <cuvs/core/export.hpp>

namespace CUVS_EXPORT cuvs {
Expand Down
2 changes: 1 addition & 1 deletion cpp/include/cuvs/preprocessing/quantize/binary.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,7 @@
#include <raft/core/host_mdarray.hpp>
#include <raft/core/host_mdspan.hpp>

#include <cuda_fp16.h>
#include <cuvs/core/cuda_fp16.hpp>
#include <cuvs/core/export.hpp>

namespace CUVS_EXPORT cuvs {
Expand Down
2 changes: 1 addition & 1 deletion cpp/include/cuvs/preprocessing/quantize/scalar.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,7 @@
#include <raft/core/host_mdarray.hpp>
#include <raft/core/host_mdspan.hpp>

#include <cuda_fp16.h>
#include <cuvs/core/cuda_fp16.hpp>
#include <cuvs/core/export.hpp>

namespace CUVS_EXPORT cuvs {
Expand Down
2 changes: 1 addition & 1 deletion cpp/include/cuvs/selection/select_k.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@

#pragma once

#include <cuda_fp16.h>
#include <cuvs/core/cuda_fp16.hpp>

#include <raft/core/device_mdspan.hpp>
#include <raft/core/resources.hpp>
Expand Down
19 changes: 13 additions & 6 deletions cpp/src/neighbors/brute_force.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,15 +1,19 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
* SPDX-License-Identifier: Apache-2.0
*/

#include <cuvs/core/cuda_fp16.hpp>

#include "./detail/knn_brute_force.cuh"

#include <cuvs/neighbors/brute_force.hpp>

#include <raft/core/copy.hpp>

namespace cuvs::neighbors::brute_force {
namespace CUVS_EXPORT cuvs {
namespace neighbors {
namespace brute_force {

template <typename T, typename DistT>
index<T, DistT>::index(raft::resources const& res)
Expand Down Expand Up @@ -227,13 +231,16 @@ void index<T, DistT>::update_dataset(
{ \
detail::search<T, int64_t, DistT, raft::col_major>( \
res, idx, queries, neighbors, distances, sample_filter); \
} \
\
template struct cuvs::neighbors::brute_force::index<T, DistT>;
}

CUVS_INST_BFKNN(float, float);
CUVS_INST_BFKNN(half, float);

template class cuvs::neighbors::brute_force::index<float, float>;
template class cuvs::neighbors::brute_force::index<half, float>;

#undef CUVS_INST_BFKNN

} // namespace cuvs::neighbors::brute_force
} // namespace brute_force
} // namespace neighbors
} // namespace CUVS_EXPORT cuvs
4 changes: 4 additions & 0 deletions cpp/src/neighbors/cagra_build_inst.cu.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,8 @@
* SPDX-License-Identifier: Apache-2.0
*/

#include <cuvs/core/cuda_fp16.hpp>

#include <cuvs/neighbors/cagra.hpp>
#include <neighbors/cagra.cuh>

Expand Down Expand Up @@ -39,4 +41,6 @@ auto build(raft::resources const& handle,
return cuvs::neighbors::cagra::build<data_t, index_t>(handle, params, dataset);
}

template struct index<data_t, index_t>;

} // namespace cuvs::neighbors::cagra
2 changes: 2 additions & 0 deletions cpp/src/neighbors/cagra_extend_inst.cu.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,8 @@
* SPDX-License-Identifier: Apache-2.0
*/

#include <cuvs/core/cuda_fp16.hpp>

#include <cuvs/neighbors/cagra.hpp>
#include <neighbors/cagra.cuh>

Expand Down
2 changes: 2 additions & 0 deletions cpp/src/neighbors/cagra_merge_inst.cu.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,8 @@
* SPDX-License-Identifier: Apache-2.0
*/

#include <cuvs/core/cuda_fp16.hpp>

#include <cuvs/neighbors/cagra.hpp>
#include <neighbors/cagra.cuh>

Expand Down
2 changes: 2 additions & 0 deletions cpp/src/neighbors/cagra_serialize_inst.cu.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,8 @@
* SPDX-License-Identifier: Apache-2.0
*/

#include <cuvs/core/cuda_fp16.hpp>

#include <neighbors/cagra_serialize.cuh>

namespace {
Expand Down
16 changes: 8 additions & 8 deletions cpp/src/neighbors/composite/index.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -111,13 +111,13 @@ void composite_index<T, IdxT, OutputIdxT>::search(
cuvs::selection::SelectAlgo::kAuto);
}

template class composite_index<float, uint32_t, uint32_t>;
template class composite_index<float, uint32_t, int64_t>;
template class composite_index<half, uint32_t, uint32_t>;
template class composite_index<half, uint32_t, int64_t>;
template class composite_index<int8_t, uint32_t, uint32_t>;
template class composite_index<int8_t, uint32_t, int64_t>;
template class composite_index<uint8_t, uint32_t, uint32_t>;
template class composite_index<uint8_t, uint32_t, int64_t>;
template class CUVS_EXPORT composite_index<float, uint32_t, uint32_t>;
template class CUVS_EXPORT composite_index<float, uint32_t, int64_t>;
template class CUVS_EXPORT composite_index<half, uint32_t, uint32_t>;
template class CUVS_EXPORT composite_index<half, uint32_t, int64_t>;
template class CUVS_EXPORT composite_index<int8_t, uint32_t, uint32_t>;
template class CUVS_EXPORT composite_index<int8_t, uint32_t, int64_t>;
template class CUVS_EXPORT composite_index<uint8_t, uint32_t, uint32_t>;
template class CUVS_EXPORT composite_index<uint8_t, uint32_t, int64_t>;

} // namespace cuvs::neighbors::composite
4 changes: 2 additions & 2 deletions cpp/src/neighbors/detail/nn_descent_gnnd.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -151,7 +151,7 @@ class BloomFilter {
};

template <typename Index_t>
struct GnndGraph {
struct CUVS_EXPORT GnndGraph {
raft::resources const& res;
static constexpr int segment_size = 32;
InternalID_t<Index_t>* h_graph;
Expand Down Expand Up @@ -192,7 +192,7 @@ struct GnndGraph {
};

template <typename Data_t = float, typename Index_t = int>
class GNND {
class CUVS_EXPORT GNND {
public:
GNND(raft::resources const& res, const BuildConfig& build_config);
GNND(const GNND&) = delete;
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/neighbors/dynamic_batching.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
* SPDX-License-Identifier: Apache-2.0
*/

Expand All @@ -21,7 +21,7 @@ namespace cuvs::neighbors::dynamic_batching {
#define CUVS_INST_DYNAMIC_BATCHING_INDEX(T, IdxT, Namespace, ...) \
template <> \
template <> \
index<T, IdxT>::index<Namespace ::__VA_ARGS__>( \
CUVS_EXPORT index<T, IdxT>::index<Namespace ::__VA_ARGS__>( \
const raft::resources& res, \
const cuvs::neighbors::dynamic_batching::index_params& params, \
const Namespace ::__VA_ARGS__& upstream_index, \
Expand Down
13 changes: 13 additions & 0 deletions cpp/src/neighbors/ivf_flat/ivf_flat_build_extend_inst.cu.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,8 @@
* SPDX-License-Identifier: Apache-2.0
*/

#include <cuvs/core/cuda_fp16.hpp>

#include <cuvs/neighbors/ivf_flat.hpp>

#include <neighbors/ivf_flat/ivf_flat_build.cuh>
Expand Down Expand Up @@ -84,3 +86,14 @@ void extend(raft::resources const& handle,
}

} // namespace cuvs::neighbors::ivf_flat

namespace cuvs::neighbors::ivf {

template void resize_list<list<cuvs::neighbors::ivf_flat::list_spec, uint32_t, data_t, index_t>>(
raft::resources const&,
std::shared_ptr<list<cuvs::neighbors::ivf_flat::list_spec, uint32_t, data_t, index_t>>&,
const list<cuvs::neighbors::ivf_flat::list_spec, uint32_t, data_t, index_t>::spec_type&,
list<cuvs::neighbors::ivf_flat::list_spec, uint32_t, data_t, index_t>::size_type,
list<cuvs::neighbors::ivf_flat::list_spec, uint32_t, data_t, index_t>::size_type);

} // namespace cuvs::neighbors::ivf
15 changes: 14 additions & 1 deletion cpp/src/neighbors/ivf_flat/ivf_flat_helpers.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,11 +1,12 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2023, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION.
* SPDX-License-Identifier: Apache-2.0
*/

#include <cstdint>

#include "../ivf_common.cuh"
#include "../ivf_list.cuh"
#include "ivf_flat_helpers.cuh"
#include <cuvs/neighbors/ivf_flat.hpp>

Expand Down Expand Up @@ -206,3 +207,15 @@ void recompute_internal_state(const raft::resources& res, index<uint8_t, int64_t
}

} // namespace cuvs::neighbors::ivf_flat::helpers

namespace cuvs::neighbors::ivf {

template CUVS_EXPORT void
resize_list<list<cuvs::neighbors::ivf_flat::list_spec, uint32_t, half, int64_t>>(

@divyegala divyegala May 21, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Was this a missing instantiation?

@vyasr vyasr May 21, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is related to your other question. nvcc + template specialization + explicit instantiation + explicit symbol visibility markup hit a number of edge cases and we had to find the magic incantation that would both be supported by nvcc and result in the right output. I won't pretend that I was particularly principled in my choice, I just tried the combinations that were closest to standards-compliant in reverse order until one produced the outcome I wanted.

raft::resources const&,
std::shared_ptr<list<cuvs::neighbors::ivf_flat::list_spec, uint32_t, half, int64_t>>&,
const list<cuvs::neighbors::ivf_flat::list_spec, uint32_t, half, int64_t>::spec_type&,
list<cuvs::neighbors::ivf_flat::list_spec, uint32_t, half, int64_t>::size_type,
list<cuvs::neighbors::ivf_flat::list_spec, uint32_t, half, int64_t>::size_type);

} // namespace cuvs::neighbors::ivf
10 changes: 5 additions & 5 deletions cpp/src/neighbors/ivf_list.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,11 +62,11 @@ list<SpecT, SizeT, SpecExtraArgs...>::list(raft::resources const& res,
}

template <typename ListT>
void resize_list(raft::resources const& res,
std::shared_ptr<ListT>& orig_list, // NOLINT
const typename ListT::spec_type& spec,
typename ListT::size_type new_used_size,
typename ListT::size_type old_used_size)
CUVS_EXPORT void resize_list(raft::resources const& res,
std::shared_ptr<ListT>& orig_list, // NOLINT
const typename ListT::spec_type& spec,
typename ListT::size_type new_used_size,
typename ListT::size_type old_used_size)
{
bool skip_resize = false;
if (orig_list) {
Expand Down
Loading
Loading