Skip to content

[Fix] Various fixes for 25.02.01 point release #695

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
merged 7 commits into from
Feb 24, 2025
Merged

[Fix] Various fixes for 25.02.01 point release #695

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
ac5cc15
[Fix] CAGRA 'merge' API compilation error under CUDA 12.4
rhdong Feb 14, 2025
8447421
fix by index::dataset_index_t
rhdong Feb 14, 2025
d4366ce
[Fix] Duplicated node issue on Cagra NN-descent
rhdong Feb 14, 2025
88b018d
Fix for a crash situation when using different threads (#680)
Feb 13, 2025
1b36e8c
Java: Fixing the broken pre-filtering for Brute Force
Feb 19, 2025
85d8644
FIX add sklearn as dependency of cuvs-bench
dantegd Feb 20, 2025
6a36092
[Fix] Brute force with prefilter java test failure
rhdong Feb 22, 2025
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
8 changes: 5 additions & 3 deletions cpp/include/cuvs/neighbors/cagra.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -337,6 +337,8 @@ struct index : cuvs::neighbors::index {
using search_params_type = cagra::search_params;
using index_type = IdxT;
using value_type = T;
using dataset_index_type = int64_t;
Copy link
Member

Choose a reason for hiding this comment

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

I think this is where @achirkin is suggesting not to hard code the type. Ideally we should use a template for this so that it can be propagated outside of this class (and not hardcoded within it).

Copy link
Contributor

Choose a reason for hiding this comment

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

I think in this case it's fine; this says "the dataset member of the cagra index uses int64_t as the indexing type", so one can argue it belongs to the index, and it's also an implementation detail of the cagra index. Before this, my problem was that it was hardcoded in two different places with no compile-time relation between those (inside the cagra index and in the merge function).


static_assert(!raft::is_narrowing_v<uint32_t, IdxT>,
"IdxT must be able to represent all values of uint32_t");

Expand Down Expand Up @@ -510,14 +512,14 @@ struct index : cuvs::neighbors::index {
*/
template <typename DatasetT>
auto update_dataset(raft::resources const& res, DatasetT&& dataset)
-> std::enable_if_t<std::is_base_of_v<cuvs::neighbors::dataset<int64_t>, DatasetT>>
-> std::enable_if_t<std::is_base_of_v<cuvs::neighbors::dataset<dataset_index_type>, DatasetT>>
{
dataset_ = std::make_unique<DatasetT>(std::move(dataset));
}

template <typename DatasetT>
auto update_dataset(raft::resources const& res, std::unique_ptr<DatasetT>&& dataset)
-> std::enable_if_t<std::is_base_of_v<neighbors::dataset<int64_t>, DatasetT>>
-> std::enable_if_t<std::is_base_of_v<neighbors::dataset<dataset_index_type>, DatasetT>>
{
dataset_ = std::move(dataset);
}
Expand Down Expand Up @@ -561,7 +563,7 @@ struct index : cuvs::neighbors::index {
cuvs::distance::DistanceType metric_;
raft::device_matrix<IdxT, int64_t, raft::row_major> graph_;
raft::device_matrix_view<const IdxT, int64_t, raft::row_major> graph_view_;
std::unique_ptr<neighbors::dataset<int64_t>> dataset_;
std::unique_ptr<neighbors::dataset<dataset_index_type>> dataset_;
};
/**
* @}
Expand Down
9 changes: 5 additions & 4 deletions cpp/src/neighbors/detail/cagra/cagra_merge.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,14 +43,16 @@ index<T, IdxT> merge(raft::resources const& handle,
const cagra::merge_params& params,
std::vector<cuvs::neighbors::cagra::index<T, IdxT>*>& indices)
{
using cagra_index_t = cuvs::neighbors::cagra::index<T, IdxT>;
using ds_idx_type = typename cagra_index_t::dataset_index_type;

std::size_t dim = 0;
std::size_t new_dataset_size = 0;
int64_t stride = -1;

for (auto index : indices) {
for (cagra_index_t* index : indices) {
RAFT_EXPECTS(index != nullptr,
"Null pointer detected in 'indices'. Ensure all elements are valid before usage.");
using ds_idx_type = decltype(index->data().n_rows());
if (auto* strided_dset = dynamic_cast<const strided_dataset<T, ds_idx_type>*>(&index->data());
strided_dset != nullptr) {
if (dim == 0) {
Expand All @@ -74,8 +76,7 @@ index<T, IdxT> merge(raft::resources const& handle,
IdxT offset = 0;

auto merge_dataset = [&](T* dst) {
for (auto index : indices) {
using ds_idx_type = decltype(index->data().n_rows());
for (cagra_index_t* index : indices) {
auto* strided_dset = dynamic_cast<const strided_dataset<T, ds_idx_type>*>(&index->data());

RAFT_CUDA_TRY(cudaMemcpy2DAsync(dst + offset * dim,
Expand Down
24 changes: 16 additions & 8 deletions cpp/src/neighbors/detail/nn_descent.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1047,24 +1047,32 @@ void GnndGraph<Index_t>::init_random_graph()
for (size_t seg_idx = 0; seg_idx < static_cast<size_t>(num_segments); seg_idx++) {
// random sequence (range: 0~nrow)
// segment_x stores neighbors which id % num_segments == x
std::vector<Index_t> rand_seq(nrow / num_segments);
std::vector<Index_t> rand_seq((nrow + num_segments - 1) / num_segments);
std::iota(rand_seq.begin(), rand_seq.end(), 0);
auto gen = std::default_random_engine{seg_idx};
std::shuffle(rand_seq.begin(), rand_seq.end(), gen);

#pragma omp parallel for
for (size_t i = 0; i < nrow; i++) {
size_t base_idx = i * node_degree + seg_idx * segment_size;
auto h_neighbor_list = h_graph + base_idx;
auto h_dist_list = h_dists.data_handle() + base_idx;
size_t base_idx = i * node_degree + seg_idx * segment_size;
auto h_neighbor_list = h_graph + base_idx;
auto h_dist_list = h_dists.data_handle() + base_idx;
size_t idx = base_idx;
size_t self_in_this_seg = 0;
for (size_t j = 0; j < static_cast<size_t>(segment_size); j++) {
size_t idx = base_idx + j;
Index_t id = rand_seq[idx % rand_seq.size()] * num_segments + seg_idx;
if ((size_t)id == i) {
id = rand_seq[(idx + segment_size) % rand_seq.size()] * num_segments + seg_idx;
idx++;
id = rand_seq[idx % rand_seq.size()] * num_segments + seg_idx;
self_in_this_seg = 1;
}
h_neighbor_list[j].id_with_flag() = id;
h_dist_list[j] = std::numeric_limits<DistData_t>::max();

h_neighbor_list[j].id_with_flag() =
j < (rand_seq.size() - self_in_this_seg) && size_t(id) < nrow
? id
: std::numeric_limits<Index_t>::max();
h_dist_list[j] = std::numeric_limits<DistData_t>::max();
idx++;
}
}
}
Expand Down
24 changes: 24 additions & 0 deletions cpp/tests/neighbors/ann_cagra.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -952,6 +952,12 @@ class AnnCagraIndexMergeTest : public ::testing::TestWithParam<AnnCagraInputs> {
(ps.k * ps.dim * 8 / 5 /*(=magic number)*/ < ps.n_rows))
GTEST_SKIP();

// Avoid splitting datasets with a size of 0
if (ps.n_rows <= 3) GTEST_SKIP();

// IVF_PQ requires the `n_rows >= n_lists`.
if (ps.n_rows < 8 && ps.build_algo == graph_build_algo::IVF_PQ) GTEST_SKIP();

size_t queries_size = ps.n_queries * ps.k;
std::vector<IdxT> indices_Cagra(queries_size);
std::vector<IdxT> indices_naive(queries_size);
Expand Down Expand Up @@ -1161,6 +1167,24 @@ inline std::vector<AnnCagraInputs> generate_inputs()
{0.995});
inputs.insert(inputs.end(), inputs2.begin(), inputs2.end());

// Corner cases for small datasets
inputs2 = raft::util::itertools::product<AnnCagraInputs>(
{2},
{3, 5, 31, 32, 64, 101},
{1, 10},
{2}, // k
{graph_build_algo::IVF_PQ, graph_build_algo::NN_DESCENT},
{search_algo::SINGLE_CTA, search_algo::MULTI_CTA, search_algo::MULTI_KERNEL},
{0}, // query size
{0},
{256},
{1},
{cuvs::distance::DistanceType::L2Expanded},
{false},
{true},
{0.995});
inputs.insert(inputs.end(), inputs2.begin(), inputs2.end());

// Varying dim and build algo.
inputs2 = raft::util::itertools::product<AnnCagraInputs>(
{100},
Expand Down
Loading