From 8be02f2afd1a6ef3fe33b539baf6ed86a9f89902 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Sat, 10 Oct 2026 22:28:55 +0000 Subject: [PATCH 1/6] [CI ONLY] Add diagnostics to IVF-Flat recall failures --- cpp/tests/neighbors/ann_ivf_flat.cuh | 344 +++++++++++++++++++++++++++ 1 file changed, 344 insertions(+) diff --git a/cpp/tests/neighbors/ann_ivf_flat.cuh b/cpp/tests/neighbors/ann_ivf_flat.cuh index 47a721bf2f..847ab77402 100644 --- a/cpp/tests/neighbors/ann_ivf_flat.cuh +++ b/cpp/tests/neighbors/ann_ivf_flat.cuh @@ -20,8 +20,15 @@ #include #include #include +#include #include +#include +#include +#include +#include +#include + namespace cuvs::neighbors::ivf_flat { struct test_ivf_sample_filter { @@ -65,6 +72,230 @@ class AnnIVFFlatTest : public ::testing::TestWithParam> { { } + // [CI ONLY] Diagnostics for the flaky recall-0 / duplicate-index failures. + void diagnose(const char* where, + const cuvs::neighbors::ivf_flat::index& idx, + const std::vector& naive_idx, + const std::vector& naive_dist, + const std::vector& ivf_idx, + const std::vector& ivf_dist, + IdxT db_offset) + { + raft::resource::sync_stream(handle_); + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + auto& os = std::cout; + const auto n_lists = idx.n_lists(); + const auto dim = idx.dim(); + const size_t n_db = ps.num_db_vecs; + const size_t k = ps.k; + const size_t nq = ps.num_queries; + os << "[DIAG " << where << "] params " << ps << " n_lists=" << n_lists << " dim=" << dim + << " veclen=" << idx.veclen() << " size=" << idx.size() + << " adaptive=" << idx.adaptive_centers() + << " conservative=" << idx.conservative_memory_allocation() << std::endl; + + // 1. Input data integrity + std::vector db_h(n_db * ps.dim), q_h(nq * ps.dim); + raft::update_host(db_h.data(), database.data(), db_h.size(), stream_); + raft::update_host(q_h.data(), search_queries.data(), q_h.size(), stream_); + raft::resource::sync_stream(handle_); + size_t db_changed = 0, q_changed = 0, db_zero = 0, q_zero = 0; + for (size_t i = 0; i < db_h.size(); i++) + db_changed += std::memcmp(&db_h[i], &db_host_[i], sizeof(DataT)) != 0; + for (size_t i = 0; i < q_h.size(); i++) + q_changed += std::memcmp(&q_h[i], &q_host_[i], sizeof(DataT)) != 0; + for (size_t i = 0; i < n_db; i++) { + bool z = true; + for (size_t d = 0; d < size_t(ps.dim); d++) + z = z && static_cast(db_h[i * ps.dim + d]) == 0.f; + db_zero += z; + } + for (size_t i = 0; i < nq; i++) { + bool z = true; + for (size_t d = 0; d < size_t(ps.dim); d++) + z = z && static_cast(q_h[i * ps.dim + d]) == 0.f; + q_zero += z; + } + os << "[DIAG " << where << "] database changed elems=" << db_changed + << " queries changed elems=" << q_changed << " zero db rows=" << db_zero + << " zero query rows=" << q_zero << std::endl; + + // 2. Centers + std::vector centers_h(size_t(n_lists) * dim); + raft::update_host(centers_h.data(), idx.centers().data_handle(), centers_h.size(), stream_); + std::vector cnorm_h; + if (idx.center_norms().has_value()) { + cnorm_h.resize(n_lists); + raft::update_host(cnorm_h.data(), idx.center_norms()->data_handle(), n_lists, stream_); + } + std::vector sizes_h(n_lists); + raft::update_host(sizes_h.data(), idx.list_sizes().data_handle(), n_lists, stream_); + std::vector inds_ptrs_h(n_lists); + raft::update_host(inds_ptrs_h.data(), idx.inds_ptrs().data_handle(), n_lists, stream_); + std::vector data_ptrs_h(n_lists); + raft::update_host(data_ptrs_h.data(), idx.data_ptrs().data_handle(), n_lists, stream_); + raft::resource::sync_stream(handle_); + size_t c_nan = 0, c_zero = 0, cn_nan = 0, cn_zero = 0, cn_mismatch = 0; + for (uint32_t l = 0; l < n_lists; l++) { + bool z = true, nan = false; + double nn = 0; + for (uint32_t d = 0; d < dim; d++) { + float v = centers_h[size_t(l) * dim + d]; + z = z && v == 0.f; + nan = nan || !std::isfinite(v); + nn += double(v) * v; + } + c_zero += z; + c_nan += nan; + if (!cnorm_h.empty()) { + cn_nan += !std::isfinite(cnorm_h[l]); + cn_zero += cnorm_h[l] == 0.f; + double expect = + ps.metric == cuvs::distance::DistanceType::CosineExpanded ? std::sqrt(nn) : nn; + cn_mismatch += std::abs(expect - cnorm_h[l]) > 1e-3 * std::max(1.0, expect); + } + } + os << "[DIAG " << where << "] centers: non-finite rows=" << c_nan << " zero rows=" << c_zero + << " center_norms: present=" << !cnorm_h.empty() << " non-finite=" << cn_nan + << " zero=" << cn_zero << " mismatch=" << cn_mismatch << std::endl; + + // 3. Lists + size_t sum = 0, empty = 0, max_sz = 0, ptr_mismatch = 0, cap_short = 0; + for (uint32_t l = 0; l < n_lists; l++) { + sum += sizes_h[l]; + empty += sizes_h[l] == 0; + max_sz = std::max(max_sz, sizes_h[l]); + auto& lst = idx.lists()[l]; + if (lst) { + if (lst->indices.data_handle() != inds_ptrs_h[l] || + lst->data.data_handle() != data_ptrs_h[l]) + ptr_mismatch++; + if (lst->indices.size() < sizes_h[l]) cap_short++; + } else if (sizes_h[l] > 0) { + ptr_mismatch++; + } + } + os << "[DIAG " << where << "] lists: sum sizes=" << sum << " empty=" << empty + << " max=" << max_sz << " ptr mismatch=" << ptr_mismatch << " capacity short=" << cap_short + << std::endl; + + std::vector seen(n_db, 0); + std::vector where_list(n_db, -1); + size_t bad_ids = 0, data_mismatch = 0, data_mismatch_lists = 0; + for (uint32_t l = 0; l < n_lists; l++) { + if (sizes_h[l] == 0 || cap_short) continue; + std::vector ids(sizes_h[l]); + raft::update_host(ids.data(), inds_ptrs_h[l], sizes_h[l], stream_); + size_t padded = raft::round_up_safe(sizes_h[l], kIndexGroupSize); + std::vector ldata(padded * dim); + raft::update_host(ldata.data(), data_ptrs_h[l], ldata.size(), stream_); + raft::resource::sync_stream(handle_); + size_t mism_here = 0; + for (uint32_t r = 0; r < sizes_h[l]; r++) { + auto id = ids[r]; + if (id < 0 || size_t(id) >= n_db) { + bad_ids++; + continue; + } + seen[id]++; + where_list[id] = l; + uint32_t vl = idx.veclen(); + for (uint32_t d = 0; d < dim; d++) { + size_t off = size_t(r / kIndexGroupSize) * kIndexGroupSize * dim + + size_t(d - d % vl) * kIndexGroupSize + size_t(r % kIndexGroupSize) * vl + + d % vl; + if (std::memcmp(&ldata[off], &db_h[size_t(id) * dim + d], sizeof(DataT)) != 0) { + mism_here++; + break; + } + } + } + data_mismatch += mism_here; + data_mismatch_lists += mism_here > 0; + } + size_t missing = 0, dup = 0; + for (size_t i = 0; i < n_db; i++) { + missing += seen[i] == 0; + dup += seen[i] > 1; + } + os << "[DIAG " << where << "] ids: out of range=" << bad_ids << " missing=" << missing + << " duplicated=" << dup << " rows whose list data != database=" << data_mismatch << " (in " + << data_mismatch_lists << " lists)" << std::endl; + + // 4. Results + size_t res_invalid = 0, res_nan = 0, res_max = 0, q_zero_match = 0; + for (size_t i = 0; i < nq; i++) { + size_t matches = 0; + for (size_t j = 0; j < k; j++) { + auto id = ivf_idx[i * k + j]; + auto dd = ivf_dist[i * k + j]; + res_invalid += id < 0 || size_t(id) >= n_db; + res_nan += std::isnan(dd); + res_max += dd == std::numeric_limits::max() || dd == std::numeric_limits::lowest() || + std::isinf(dd); + for (size_t jj = 0; jj < k; jj++) + if (std::abs(dd - naive_dist[i * k + jj]) <= 0.001 || id == naive_idx[i * k + jj]) { + matches++; + break; + } + } + q_zero_match += matches == 0; + } + os << "[DIAG " << where << "] results: invalid ids=" << res_invalid << " NaN dists=" << res_nan + << " inf/max dists=" << res_max << " queries with 0 matches=" << q_zero_match << std::endl; + + // 5. Per-query details for the first queries + auto qnorm = [&](size_t q) { + double s = 0; + for (uint32_t d = 0; d < dim; d++) { + double v = static_cast(q_h[q * dim + d]); + s += v * v; + } + return std::sqrt(s); + }; + for (size_t q = 0; q < std::min(3, nq); q++) { + os << "[DIAG " << where << "] query " << q << " = ("; + for (uint32_t d = 0; d < std::min(dim, 8); d++) + os << static_cast(q_h[q * dim + d]) << (d + 1 < dim ? "," : ""); + os << ")" << std::endl << " naive: "; + for (size_t j = 0; j < k; j++) { + auto id = naive_idx[q * k + j] + db_offset; + os << id << ":" << naive_dist[q * k + j] << "[L" + << (size_t(id) < n_db ? where_list[id] : -2) << "] "; + } + os << std::endl << " ivf: "; + for (size_t j = 0; j < k; j++) { + auto id = ivf_idx[q * k + j]; + os << id << ":" << ivf_dist[q * k + j] << "[L" + << (id >= 0 && size_t(id) < n_db ? where_list[id] : -2) << "] "; + } + os << std::endl; + // coarse ranking (host), cosine/inner/L2 approximations + std::vector> cd(n_lists); + double qn = qnorm(q); + for (uint32_t l = 0; l < n_lists; l++) { + double dot = 0, l2 = 0; + for (uint32_t d = 0; d < dim; d++) { + double c = centers_h[size_t(l) * dim + d]; + double x = static_cast(q_h[q * dim + d]); + dot += c * x; + l2 += (c - x) * (c - x); + } + double cn = cnorm_h.empty() ? 1.0 : cnorm_h[l]; + double v = ps.metric == cuvs::distance::DistanceType::CosineExpanded ? -dot / (qn * cn) + : ps.metric == cuvs::distance::DistanceType::InnerProduct ? -dot + : l2; + if (std::isnan(v)) v = 1e30; + cd[l] = {v, l}; + } + std::sort(cd.begin(), cd.end()); + os << " coarse top-" << ps.nprobe << " (list:dist:size): "; + for (size_t j = 0; j < std::min(ps.nprobe, n_lists); j++) + os << cd[j].second << ":" << cd[j].first << ":" << sizes_h[cd[j].second] << " "; + os << std::endl; + } + } + void testIVFFlat() { size_t queries_size = ps.num_queries * ps.k; @@ -204,6 +435,33 @@ class AnnIVFFlatTest : public ::testing::TestWithParam> { indices_ivfflat.data(), indices_ivfflat_dev.data(), queries_size, stream_); raft::resource::sync_stream(handle_); + { + float diag_eps = std::is_same_v ? 0.005 : 0.001; + if (!eval_neighbours(indices_naive, + indices_ivfflat, + distances_naive, + distances_ivfflat, + ps.num_queries, + ps.k, + diag_eps, + min_recall)) { + diagnose("testIVFFlat/index_loaded", + index_loaded, + indices_naive, + distances_naive, + indices_ivfflat, + distances_ivfflat, + 0); + diagnose("testIVFFlat/index_2", + index_2, + indices_naive, + distances_naive, + indices_ivfflat, + distances_ivfflat, + 0); + } + } + // Test the centroid invariants if (index_2.adaptive_centers()) { // The centers must be up-to-date with the corresponding data @@ -474,6 +732,85 @@ class AnnIVFFlatTest : public ::testing::TestWithParam> { raft::update_host( indices_ivfflat.data(), indices_ivfflat_dev.data_handle(), queries_size, stream_); raft::resource::sync_stream(handle_); + + float diag_eps = std::is_same_v ? 0.005 : 0.001; + if (!eval_neighbours(indices_naive, + indices_ivfflat, + distances_naive, + distances_ivfflat, + ps.num_queries, + ps.k, + diag_eps, + min_recall)) { + diagnose("testFilter", + index, + indices_naive, + distances_naive, + indices_ivfflat, + distances_ivfflat, + 0); + // Repeat the same search on the same index to tell a corrupted index from a transient + // search-time error. + std::vector idx2(queries_size); + std::vector dist2(queries_size); + ivf_flat::search(handle_, + search_params, + index, + search_queries_view, + indices_ivfflat_dev.view(), + distances_ivfflat_dev.view(), + bitset_filter_obj); + raft::update_host( + dist2.data(), distances_ivfflat_dev.data_handle(), queries_size, stream_); + raft::update_host(idx2.data(), indices_ivfflat_dev.data_handle(), queries_size, stream_); + raft::resource::sync_stream(handle_); + std::cout << "[DIAG testFilter] re-search with the same filter:" << std::endl; + (void)eval_neighbours(indices_naive, + idx2, + distances_naive, + dist2, + ps.num_queries, + ps.k, + diag_eps, + min_recall); + // Rebuild the bitset and search again. + cuvs::core::bitset bitset2( + handle_, removed_indices.view(), ps.num_db_vecs); + std::vector bits_h(bitset2.view().n_elements()); + std::vector bits_old(removed_indices_bitset.view().n_elements()); + raft::update_host(bits_h.data(), bitset2.view().data(), bits_h.size(), stream_); + raft::update_host( + bits_old.data(), removed_indices_bitset.view().data(), bits_old.size(), stream_); + raft::resource::sync_stream(handle_); + size_t diff_words = 0, zero_bits = 0; + for (size_t i = 0; i < bits_h.size(); i++) { + diff_words += bits_h[i] != bits_old[i]; + zero_bits += 32 - __builtin_popcount(bits_old[i]); + } + std::cout << "[DIAG testFilter] bitset words=" << bits_old.size() + << " differing words vs fresh bitset=" << diff_words + << " zero bits in original=" << zero_bits << std::endl; + ivf_flat::search(handle_, + search_params, + index, + search_queries_view, + indices_ivfflat_dev.view(), + distances_ivfflat_dev.view(), + cuvs::neighbors::filtering::bitset_filter(bitset2.view())); + raft::update_host( + dist2.data(), distances_ivfflat_dev.data_handle(), queries_size, stream_); + raft::update_host(idx2.data(), indices_ivfflat_dev.data_handle(), queries_size, stream_); + raft::resource::sync_stream(handle_); + std::cout << "[DIAG testFilter] re-search with a fresh bitset:" << std::endl; + (void)eval_neighbours(indices_naive, + idx2, + distances_naive, + dist2, + ps.num_queries, + ps.k, + diag_eps, + min_recall); + } } float eps = std::is_same_v ? 0.005 : 0.001; ASSERT_TRUE(eval_neighbours(indices_naive, @@ -505,6 +842,11 @@ class AnnIVFFlatTest : public ::testing::TestWithParam> { handle_, r, search_queries.data(), ps.num_queries * ps.dim, DataT(1), DataT(20)); } raft::resource::sync_stream(handle_); + db_host_.resize(ps.num_db_vecs * ps.dim); + q_host_.resize(ps.num_queries * ps.dim); + raft::update_host(db_host_.data(), database.data(), db_host_.size(), stream_); + raft::update_host(q_host_.data(), search_queries.data(), q_host_.size(), stream_); + raft::resource::sync_stream(handle_); } void TearDown() override @@ -520,6 +862,8 @@ class AnnIVFFlatTest : public ::testing::TestWithParam> { AnnIvfFlatInputs ps; rmm::device_uvector database; rmm::device_uvector search_queries; + std::vector db_host_; + std::vector q_host_; }; const std::vector> inputs = { From 8ed8cfff7bbad9362c40af494dcdb41b348401f2 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Sat, 10 Oct 2026 22:29:15 +0000 Subject: [PATCH 2/6] [CI ONLY] Repeat IVF-Flat int8/uint8 cosine test on V100 and H100 --- .github/workflows/pr.yaml | 723 +------------------------------------- ci/test_cpp.sh | 39 +- 2 files changed, 26 insertions(+), 736 deletions(-) diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index a6438c1a8d..c4a2d23643 100644 --- a/.github/workflows/pr.yaml +++ b/.github/workflows/pr.yaml @@ -11,450 +11,18 @@ jobs: pr-builder: needs: - build-details - - check-nightly-ci - - changed-files - - checks - - check-c-abi - - security-suite - conda-cpp-build - conda-cpp-tests - - conda-cpp-checks - - conda-python-build - - conda-python-tests - - rocky8-clib-standalone-build - - rocky8-clib-tests - - conda-java-build - - conda-java-tests - - conda-lucene-build - - conda-lucene-tests - - rust-build - - go-build - - docs-build - - docs-build-matrix - - wheel-build-libcuvs - - wheel-build-cuvs - - wheel-tests-cuvs - - devcontainer - - telemetry-setup uses: rapidsai/shared-workflows/.github/workflows/pr-builder.yaml@main permissions: contents: read if: always() with: needs: ${{ toJSON(needs) }} - telemetry-setup: - continue-on-error: true - runs-on: ubuntu-latest - permissions: - contents: read - env: - OTEL_SERVICE_NAME: 'pr-cuvs' - steps: - - name: Telemetry setup - if: ${{ vars.TELEMETRY_ENABLED == 'true' }} - uses: rapidsai/shared-actions/telemetry-dispatch-stash-base-env-vars@main build-details: uses: rapidsai/shared-workflows/.github/workflows/compute-build-details.yaml@main - check-nightly-ci: - runs-on: ubuntu-latest - permissions: - actions: read - contents: read - id-token: write - pull-requests: read - env: - GH_TOKEN: ${{ github.token }} - steps: - - name: Get PR Info - id: get-pr-info - uses: nv-gha-runners/get-pr-info@main - - name: Check if nightly CI is passing - uses: rapidsai/shared-actions/check_nightly_success/dispatch@main - with: - repo: ${{ github.repository }} - target-branch: ${{ fromJSON(steps.get-pr-info.outputs.pr-info).base.ref }} - max-days-without-success: 30 - changed-files: - needs: telemetry-setup - permissions: - actions: read - contents: read - packages: read - pull-requests: read - uses: rapidsai/shared-workflows/.github/workflows/changed-files.yaml@main - with: - files_yaml: | - build_docs: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!Dockerfile' - - '!SECURITY.md' - - '!ci/build_wheel*.sh' - - '!ci/check_style.sh' - - '!ci/release/update-version.sh' - - '!ci/run*.sh' - - '!ci/test*.sh' - - '!ci/validate_wheel.sh' - - '!thirdparty/LICENSES/**' - test_cpp: - - '.github/workflows/pr.yaml' - - 'RAPIDS_BRANCH' - - 'VERSION' - - 'build.sh' - - 'c/**' - - 'ci/build_cpp.sh' - - 'ci/get_test_data.sh' - - 'ci/test_cpp.sh' - - 'cmake/**' - - 'conda/recipes/libcuvs/**' - - 'cpp/**' - - 'dependencies.yaml' - - '!**/*/REVIEW_GUIDELINES.md' - - '!cpp/doxygen/**' - test_standalone_c: - - '.github/workflows/pr.yaml' - - 'RAPIDS_BRANCH' - - 'VERSION' - - 'build.sh' - - 'c/**' - - 'ci/standalone_c/**' - - 'ci/test_standalone_c.sh' - - 'cmake/**' - - 'cpp/**' - - 'dependencies.yaml' - - '!**/*/REVIEW_GUIDELINES.md' - - '!cpp/doxygen/**' - test_java: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!CODEOWNERS' - - '!Dockerfile' - - '!README.md' - - '!SECURITY.md' - - '!ci/build_docs.sh' - - '!ci/build_go.sh' - - '!ci/build_python.sh' - - '!ci/build_rust.sh' - - '!ci/build_wheel*.sh' - - '!ci/checkstyle.sh' - - '!ci/run_cuvs_pytests.sh' - - '!ci/test_cpp.sh' - - '!ci/test_python.sh' - - '!ci/test_wheel_cuvs.sh' - - '!ci/validate_wheel.sh' - - '!conda/environments/**' - - '!conda/recipes/cuvs/**' - - '!conda/recipes/cuvs-bench/**' - - '!conda/recipes/cuvs-bench-cpu/**' - - '!docs/**' - - '!examples/.clang-format' - - '!examples/.clang-tidy' - - '!examples/.clangd' - - '!examples/README.md' - - '!examples/build.sh' - - '!examples/c/**' - - '!examples/cmake/**' - - '!examples/cpp/**' - - '!examples/go/**' - - '!fern/**' - - '!go/**' - - '!img/**' - - '!notebooks/**' - - '!pyproject.toml' - - '!python/**' - - '!rust/**' - - '!thirdparty/LICENSES/**' - test_python_conda: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!ci/build_go.sh' - - '!ci/build_java.sh' - - '!ci/build_lucene.sh' - - '!ci/build_rust.sh' - - '!ci/check_style.sh' - - '!ci/release/update-version.sh' - - '!ci/run_ctests.sh' - - '!ci/test_cpp.sh' - - '!ci/test_java.sh' - - '!ci/test_lucene.sh' - - '!ci/test_standalone_c.sh' - - '!ci/test_wheel_cuvs.sh' - - '!ci/validate_wheel.sh' - - '!cpp/.clang-format' - - '!cpp/.clang-tidy' - - '!CODEOWNERS' - - '!Dockerfile' - - '!README.md' - - '!SECURITY.md' - - '!docs/**' - - '!examples/**' - - '!fern/**' - - '!go/**' - - '!img/**' - - '!java/**' - - '!notebooks/**' - - '!rust/**' - - '!thirdparty/LICENSES/**' - test_python_wheels: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!ci/build_cpp.sh' - - '!ci/build_docs.sh' - - '!ci/build_go.sh' - - '!ci/build_java.sh' - - '!ci/build_lucene.sh' - - '!ci/build_python.sh' - - '!ci/build_rust.sh' - - '!ci/build_standalone_c.sh' - - '!ci/check_style.sh' - - '!ci/release/update-version.sh' - - '!ci/run_ctests.sh' - - '!ci/test_cpp.sh' - - '!ci/test_java.sh' - - '!ci/test_lucene.sh' - - '!ci/test_python.sh' - - '!ci/test_standalone_c.sh' - - '!ci/validate_wheel.sh' - - '!CODEOWNERS' - - '!Dockerfile' - - '!README.md' - - '!SECURITY.md' - - '!conda/**' - - '!cpp/.clang-format' - - '!cpp/.clang-tidy' - - '!docs/**' - - '!examples/**' - - '!fern/**' - - '!java/**' - - '!go/**' - - '!img/**' - - '!notebooks/**' - - '!rust/**' - - '!thirdparty/LICENSES/**' - test_rust: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!CODEOWNERS' - - '!Dockerfile' - - '!README.md' - - '!SECURITY.md' - - '!ci/build_docs.sh' - - '!ci/build_go.sh' - - '!ci/build_java.sh' - - '!ci/build_lucene.sh' - - '!ci/build_python.sh' - - '!ci/build_wheel*.sh' - - '!ci/check_style.sh' - - '!ci/release/update-version.sh' - - '!ci/run_ctests.sh' - - '!ci/run_cuvs_pytests.sh' - - '!ci/test_cpp.sh' - - '!ci/test_java.sh' - - '!ci/test_lucene.sh' - - '!ci/test_python.sh' - - '!ci/test_wheel_cuvs.sh' - - '!ci/validate_wheel.sh' - - '!docs/**' - - '!examples/**' - - '!fern/**' - - '!go/**' - - '!img/**' - - '!java/**' - - '!notebooks/**' - - '!python/**' - - '!thirdparty/LICENSES/**' - test_go: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!CODEOWNERS' - - '!Dockerfile' - - '!README.md' - - '!SECURITY.md' - - '!ci/build_docs.sh' - - '!ci/build_java.sh' - - '!ci/build_lucene.sh' - - '!ci/build_python.sh' - - '!ci/build_rust.sh' - - '!ci/build_wheel*.sh' - - '!ci/check_style.sh' - - '!ci/release/update-version.sh' - - '!ci/run_ctests.sh' - - '!ci/run_cuvs_pytests.sh' - - '!ci/test_cpp.sh' - - '!ci/test_java.sh' - - '!ci/test_lucene.sh' - - '!ci/test_python.sh' - - '!ci/test_standalone_c.sh' - - '!ci/test_wheel_cuvs.sh' - - '!ci/validate_wheel.sh' - - '!docs/**' - - '!examples/.clang-format' - - '!examples/.clang-tidy' - - '!examples/.clangd' - - '!examples/README.md' - - '!examples/build.sh' - - '!examples/c/**' - - '!examples/cmake/**' - - '!examples/cpp/**' - - '!examples/java/**' - - '!fern/**' - - '!img/**' - - '!java/**' - - '!notebooks/**' - - '!python/**' - - '!rust/**' - - '!thirdparty/LICENSES/**' - checks: - needs: telemetry-setup - uses: rapidsai/shared-workflows/.github/workflows/checks.yaml@main - permissions: - contents: read - with: - enable_check_generated_files: false - ignored_pr_jobs: "telemetry-summarize" - check-c-abi: - needs: telemetry-setup - permissions: - actions: read - contents: read - issues: write - pull-requests: read - uses: ./.github/workflows/check-c-abi.yaml - security-suite: - permissions: - actions: read - contents: read - id-token: write # OIDC -> Vault -> nvcr.io image pull - security-events: write # publish redacted SARIF to code scanning - uses: ./.github/workflows/security-suite.yml conda-cpp-build: - needs: [build-details, checks] + needs: [build-details] permissions: actions: read contents: read @@ -468,8 +36,9 @@ jobs: build-datetime: ${{ needs.build-details.outputs.build-datetime }} node_type: cpu16 script: ci/build_cpp.sh + matrix_filter: map(select(.ARCH == "amd64" and .CUDA_VER == "12.9.2")) conda-cpp-tests: - needs: [conda-cpp-build, changed-files] + needs: [conda-cpp-build] permissions: actions: read contents: read @@ -478,16 +47,10 @@ jobs: pull-requests: read secrets: inherit # zizmor: ignore[secrets-inherit] uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@main - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_cpp - # Split the C++ tests across several jobs for each CUDA/GPU configuration - strategy: - fail-fast: false - matrix: - shard: [1, 2, 3, 4] - num_shards: [4] with: build_type: pull-request - script: "ci/test_cpp.sh ${{ matrix.shard }} ${{ matrix.num_shards }}" + script: "ci/test_cpp.sh" + matrix_filter: 'map(select(.ARCH == "amd64" and (.GPU == "v100" or .GPU == "h100"))) | [.[] | if .GPU == "v100" then (., . + {COPY: 1}, . + {COPY: 2}) else . end]' cache-paths: .cache/cuda-jit cache-key-prefix: libcuvs-cudajit-v1 # Caches are immutable. Re-key when the JIT-relevant sources change so a completed run can @@ -518,279 +81,3 @@ jobs: cache-read-only: true cache-environment: | CUDA_CACHE_PATH=.cache/cuda-jit - conda-cpp-checks: - needs: conda-cpp-build - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-post-build-checks.yaml@main - with: - build_type: pull-request - symbol_exclusions: (void (thrust::|cub::)) - package_name: libcuvs - conda-python-build: - needs: [build-details, conda-cpp-build] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@main - with: - build_type: pull-request - build-datetime: ${{ needs.build-details.outputs.build-datetime }} - script: ci/build_python.sh - # Build a conda package for each CUDA x ARCH x minimum supported Python version - matrix_filter: group_by({CUDA_VER, ARCH}) | map(min_by(.PY_VER | split(".") | map(tonumber))) - conda-python-tests: - needs: [conda-python-build, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@main - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_conda - with: - build_type: pull-request - script: ci/test_python.sh - rocky8-clib-standalone-build: - needs: checks - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/clib-build.yaml - with: - build_type: pull-request - rocky8-clib-tests: - needs: [rocky8-clib-standalone-build, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_standalone_c - with: - build_type: pull-request - container-image-repository: "rapidsai/citestwheel" - script: "ci/test_standalone_c.sh" - conda-java-build: - needs: [conda-cpp-build, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_java || fromJSON(needs.changed-files.outputs.changed_file_groups).test_cpp - # Artifacts are not published from these jobs, so it's safe to run for multiple CUDA versions. - # If these jobs start producing artifacts, the names will have to differentiate between CUDA versions. - with: - build_type: pull-request - node_type: cpu8 - matrix_filter: map(select(.ARCH == "amd64")) - script: "ci/build_java.sh" - artifact-name-prefix: "cuvs-java" - file_to_upload: "java/cuvs-java/target/" - conda-java-tests: - needs: conda-java-build - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - with: - build_type: pull-request - script: "ci/test_java.sh" - conda-lucene-build: - # Depends on the Java job for the cuvs-java artifact that cuvs-lucene builds against. - needs: conda-java-build - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - # Artifacts are not published from these jobs, so it's safe to run for multiple CUDA versions. - # If these jobs start producing artifacts, the names will have to differentiate between CUDA versions. - with: - build_type: pull-request - node_type: cpu8 - matrix_filter: map(select(.ARCH == "amd64")) - script: "ci/build_lucene.sh" - artifact-name-prefix: "cuvs-lucene" - file_to_upload: "java/cuvs-lucene/target/" - conda-lucene-tests: - needs: conda-lucene-build - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - with: - build_type: pull-request - script: "ci/test_lucene.sh" - rust-build: - needs: [conda-cpp-build, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_rust || fromJSON(needs.changed-files.outputs.changed_file_groups).test_cpp - # Artifacts are not published from these jobs, so it's safe to run for multiple CUDA versions. - # If these jobs start producing artifacts, the names will have to differentiate between CUDA versions. - with: - build_type: pull-request - matrix_filter: map(select(.ARCH == "amd64")) - script: "ci/build_rust.sh" - go-build: - needs: [conda-cpp-build, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_go || fromJSON(needs.changed-files.outputs.changed_file_groups).test_cpp - # Artifacts are not published from these jobs, so it's safe to run for multiple CUDA versions. - # If these jobs start producing artifacts, the names will have to differentiate between CUDA versions. - with: - build_type: pull-request - matrix_filter: map(select(.ARCH == "amd64")) - script: "ci/build_go.sh" - docs-build-matrix: - permissions: - contents: read - uses: rapidsai/shared-workflows/.github/workflows/compute-matrix.yaml@main - with: - build_type: pull-request - matrix_name: docs-build - docs-build: - needs: [docs-build-matrix] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main - if: fromJSON(needs.changed-files.outputs.changed_file_groups).build_docs - strategy: - fail-fast: false - matrix: ${{ fromJSON(needs.docs-build-matrix.outputs.matrix) }} - with: - build_type: pull-request - node_type: "gpu-${{ matrix.GPU }}-${{ matrix.DRIVER }}-1" - arch: "${{ matrix.arch }}" - container_image: "rapidsai/ci-conda:26.12-cuda${{ matrix.CUDA_VER }}-${{ matrix.LINUX_VER }}-py${{ matrix.PY_VER }}" - script: "ci/build_docs.sh" - wheel-build-libcuvs: - needs: [build-details, checks] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main - with: - build_type: pull-request - build-datetime: ${{ needs.build-details.outputs.build-datetime }} - node_type: cpu16 - script: ci/build_wheel_libcuvs.sh - # build for every combination of arch and CUDA version, but only for the latest Python - matrix_filter: group_by([.ARCH, (.CUDA_VER|split(".")|map(tonumber)|.[0])]) | map(max_by(.PY_VER|split(".")|map(tonumber))) - package-name: libcuvs - package-type: cpp - wheel-build-cuvs: - needs: [build-details, wheel-build-libcuvs] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main - with: - build_type: pull-request - build-datetime: ${{ needs.build-details.outputs.build-datetime }} - node_type: cpu8 - script: ci/build_wheel_cuvs.sh - package-name: cuvs - package-type: python - # Build a wheel for each CUDA x ARCH x minimum supported Python version - matrix_filter: group_by({CUDA_VER, ARCH}) | map(min_by(.PY_VER | split(".") | map(tonumber))) - wheel-tests-cuvs: - needs: [wheel-build-cuvs, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@main - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_wheels - with: - build_type: pull-request - script: ci/test_wheel_cuvs.sh - devcontainer: - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - needs: telemetry-setup - uses: rapidsai/shared-workflows/.github/workflows/build-in-devcontainer.yaml@main - with: - arch: '["amd64", "arm64"]' - cuda: '["13.3"]' - node_type: "cpu8" - env: | - SCCACHE_DIST_MAX_RETRIES=inf - SCCACHE_SERVER_LOG=sccache=debug - SCCACHE_DIST_FALLBACK_TO_LOCAL_COMPILE=false - build_command: | - sccache --zero-stats; - build-all -j0 --verbose 2>&1 | tee telemetry-artifacts/build.log; - sccache --show-adv-stats | tee telemetry-artifacts/sccache-stats.txt; - telemetry-summarize: - # This job must use a self-hosted runner to record telemetry traces. - runs-on: linux-amd64-cpu4 - needs: pr-builder - permissions: - actions: write - if: ${{ vars.TELEMETRY_ENABLED == 'true' && !cancelled() }} - continue-on-error: true - steps: - - name: Telemetry summarize - uses: rapidsai/shared-actions/telemetry-dispatch-summarize@main - env: - GH_TOKEN: ${{ github.token }} diff --git a/ci/test_cpp.sh b/ci/test_cpp.sh index 0e97a0ddef..5fe18ea172 100755 --- a/ci/test_cpp.sh +++ b/ci/test_cpp.sh @@ -4,11 +4,6 @@ set -euo pipefail -# Usage: ci/test_cpp.sh [SHARD NUM_SHARDS] -# With SHARD and NUM_SHARDS (1 <= SHARD <= NUM_SHARDS), only every NUM_SHARDS-th libcuvs test is -# run, starting from test number SHARD, so the tests can be split across several CI jobs. -SHARD=${1:-1} -NUM_SHARDS=${2:-1} . /opt/conda/etc/profile.d/conda.sh @@ -52,23 +47,31 @@ nvidia-smi # RAPIDS_DATASET_ROOT_DIR is used by test scripts RAPIDS_DATASET_ROOT_DIR=${RAPIDS_TESTS_DIR}/dataset export RAPIDS_DATASET_ROOT_DIR -./ci/get_test_data.sh --NEIGHBORS_ANN_VAMANA_TEST - EXITCODE=0 trap "EXITCODE=1" ERR set +e -# Run Python build utilities tests (once, in the first shard) -if [[ "${SHARD}" == "1" ]]; then - rapids-logger "Run libcuvs Python build utilities tests" - pytest cpp/tests/python -fi - -# Run libcuvs gtests from libcuvs-tests package -rapids-logger "Run libcuvs tests (shard ${SHARD} of ${NUM_SHARDS})" -pushd "$CONDA_PREFIX"/bin/gtests/libcuvs -timeout -v --signal=SIGINT --kill-after=60s 100m ctest -j8 --output-on-failure -I "${SHARD},,${NUM_SHARDS}" -popd +# [CI ONLY] Repeat the flaky IVF-Flat int8/uint8 cosine test. +BIN="$CONDA_PREFIX"/bin/gtests/libcuvs/NEIGHBORS_ANN_IVF_FLAT_TEST +repeat_test() { + local name=$1 filter=$2 repeat=$3 + local log="${RAPIDS_TESTS_DIR}/${name}.log" + rapids-logger "Repeat --gtest_filter=${filter} ${repeat} times" + timeout -v --signal=SIGINT --kill-after=60s 150m "${BIN}" --gtest_filter="${filter}" --gtest_repeat="${repeat}" > "${log}" 2>&1 + echo "Exit status: $?" + tail -n 5 "${log}" + echo "Failure details:" + grep -E -B3 -A6 "Failure$" "${log}" | head -n 400 || true + echo "Diagnostics:" + grep -E -A0 "DIAG|^ (naive|ivf|coarse)|Recall = 0\.[0-8]|Duplicated" "${log}" | head -n 2000 || true + echo "Failures:" + grep -E "^\[ FAILED \] .* \([0-9]+ ms\)$" "${log}" | sed -E 's/, where GetParam.*//; s/ \([0-9]+ ms\)//' | sort | uniq -c || true + echo "Passes:" + grep -E "^\[ OK \]" "${log}" | sed -E 's/ \([0-9]+ ms\)//' | sort | uniq -c || true +} +nvidia-smi -q | grep -E "Driver Version|CUDA Version|Product Name" || true +repeat_test cosine2 '*AnnIVFFlatTestF_int8.AnnIVFFlat/2:*AnnIVFFlatTestF_uint8.AnnIVFFlat/2' 800 +repeat_test int8_uint8_all '*AnnIVFFlatTestF_int8*:*AnnIVFFlatTestF_uint8*' 4 rapids-logger "Test script exiting with value: $EXITCODE" exit ${EXITCODE} From 04fbec451396bb80e5df430067ee34398f9960a5 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Sat, 10 Oct 2026 22:36:22 +0000 Subject: [PATCH 3/6] Test IVF-Flat cosine search with zero-norm centers --- .../ann_ivf_flat/test_float_int64_t.cu | 78 +++++++++++++++++++ 1 file changed, 78 insertions(+) diff --git a/cpp/tests/neighbors/ann_ivf_flat/test_float_int64_t.cu b/cpp/tests/neighbors/ann_ivf_flat/test_float_int64_t.cu index fee22db3c5..94916812fe 100644 --- a/cpp/tests/neighbors/ann_ivf_flat/test_float_int64_t.cu +++ b/cpp/tests/neighbors/ann_ivf_flat/test_float_int64_t.cu @@ -10,10 +10,13 @@ #include #include #include +#include +#include #include #include #include +#include namespace cuvs::neighbors::ivf_flat { @@ -48,6 +51,81 @@ TEST(IVFFlatSerialization, StreamRoundTrip) EXPECT_EQ(loaded.n_lists(), index.n_lists()); } +// k-means can leave a cluster empty, which gives it a zero center. With the cosine metric, the +// coarse search must not produce NaN distances for such lists, or the probed lists are wrong. +TEST(AnnIVFFlatTest, CosineZeroNormCenters) +{ + raft::resources handle; + auto stream = raft::resource::get_cuda_stream(handle); + + constexpr int64_t n_rows = 10000; + constexpr int64_t n_queries = 500; + constexpr int64_t dim = 8; + constexpr uint32_t n_lists = 1024; + constexpr uint32_t k = 16; + + auto data = raft::make_device_matrix(handle, n_rows, dim); + auto queries = raft::make_device_matrix(handle, n_queries, dim); + raft::random::RngState rng(1234ULL); + raft::random::uniform(handle, rng, data.data_handle(), data.size(), 0.1f, 2.0f); + raft::random::uniform(handle, rng, queries.data_handle(), queries.size(), 0.1f, 2.0f); + + index_params params; + params.n_lists = n_lists; + params.metric = cuvs::distance::DistanceType::CosineExpanded; + params.add_data_on_build = false; + auto idx = build(handle, params, raft::make_const_mdspan(data.view())); + + // Zero the centers of every other list, as if k-means had left them empty. + std::vector centers_h(idx.centers().size()); + raft::update_host(centers_h.data(), idx.centers().data_handle(), centers_h.size(), stream); + raft::resource::sync_stream(handle); + for (uint32_t l = 0; l < n_lists; l += 2) { + std::fill_n(centers_h.begin() + l * dim, dim, 0.0f); + } + raft::update_device(idx.centers().data_handle(), centers_h.data(), centers_h.size(), stream); + extend(handle, raft::make_const_mdspan(data.view()), std::nullopt, &idx); + + // Probe exactly as many lists as there are non-zero centers. + search_params sparams; + sparams.n_probes = n_lists / 2; + auto neighbors = raft::make_device_matrix(handle, n_queries, k); + auto distances = raft::make_device_matrix(handle, n_queries, k); + search(handle, + sparams, + idx, + raft::make_const_mdspan(queries.view()), + neighbors.view(), + distances.view()); + + auto naive_neighbors = raft::make_device_matrix(handle, n_queries, k); + auto naive_distances = raft::make_device_matrix(handle, n_queries, k); + cuvs::neighbors::naive_knn(handle, + naive_distances.data_handle(), + naive_neighbors.data_handle(), + queries.data_handle(), + data.data_handle(), + n_queries, + n_rows, + dim, + k, + params.metric); + + std::vector neighbors_h(n_queries * k), naive_neighbors_h(n_queries * k); + std::vector distances_h(n_queries * k), naive_distances_h(n_queries * k); + raft::update_host(neighbors_h.data(), neighbors.data_handle(), neighbors_h.size(), stream); + raft::update_host(distances_h.data(), distances.data_handle(), distances_h.size(), stream); + raft::update_host( + naive_neighbors_h.data(), naive_neighbors.data_handle(), naive_neighbors_h.size(), stream); + raft::update_host( + naive_distances_h.data(), naive_distances.data_handle(), naive_distances_h.size(), stream); + raft::resource::sync_stream(handle); + + // Every data row is in a list with a non-zero center, and all of those lists are probed. + EXPECT_TRUE(eval_neighbours( + naive_neighbors_h, neighbors_h, naive_distances_h, distances_h, n_queries, k, 0.001, 0.99)); +} + TEST(AnnIVFFlatTest, RepeatedExtendCopyPreservesSharedListWithinCapacity) { raft::resources handle; From 8bd6e350bbab71ccf45b780e39cf87bd25794506 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Sat, 10 Oct 2026 22:36:26 +0000 Subject: [PATCH 4/6] Guard IVF-Flat cosine coarse search against zero-norm centers --- cpp/src/neighbors/ivf_flat/ivf_flat_search.cuh | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/cpp/src/neighbors/ivf_flat/ivf_flat_search.cuh b/cpp/src/neighbors/ivf_flat/ivf_flat_search.cuh index 6117935712..4d2d85a8f5 100644 --- a/cpp/src/neighbors/ivf_flat/ivf_flat_search.cuh +++ b/cpp/src/neighbors/ivf_flat/ivf_flat_search.cuh @@ -171,7 +171,11 @@ void search_impl(raft::resources const& handle, [=] __device__(const uint32_t idx, const float dist) { const auto query = idx / n_lists; const auto cluster = idx % n_lists; - return dist / (q_norm_ptr[query] * index_center_norm_ptr[cluster]); + // k-means leaves a zero center for a list that ends up empty. Cosine distance involving a + // zero-norm vector is defined as 1 (similarity 0) instead of NaN, which would corrupt the + // selection of the lists to probe. + const float denom = q_norm_ptr[query] * index_center_norm_ptr[cluster]; + return denom > 0.0f ? dist / denom : 0.0f; }, raft::make_const_mdspan(distance_buffer_dev_view)); } From 2df4311e6fb8eb871e191135a5e4eb2f8390ad5b Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Sun, 11 Oct 2026 03:12:24 +0000 Subject: [PATCH 5/6] Make the zero-norm center test pad the probes with list 0 --- .../ann_ivf_flat/test_float_int64_t.cu | 30 ++++++++++--------- 1 file changed, 16 insertions(+), 14 deletions(-) diff --git a/cpp/tests/neighbors/ann_ivf_flat/test_float_int64_t.cu b/cpp/tests/neighbors/ann_ivf_flat/test_float_int64_t.cu index 94916812fe..a70d41a09f 100644 --- a/cpp/tests/neighbors/ann_ivf_flat/test_float_int64_t.cu +++ b/cpp/tests/neighbors/ann_ivf_flat/test_float_int64_t.cu @@ -51,18 +51,22 @@ TEST(IVFFlatSerialization, StreamRoundTrip) EXPECT_EQ(loaded.n_lists(), index.n_lists()); } -// k-means can leave a cluster empty, which gives it a zero center. With the cosine metric, the -// coarse search must not produce NaN distances for such lists, or the probed lists are wrong. +// k-means gives an empty cluster a zero center. With the cosine metric, the coarse search must +// still rank such lists instead of producing NaN distances: when fewer than n_probes lists have a +// finite distance, the NaN columns are never selected and the remaining probe slots are padded +// with list 0, which then is scanned several times and returns duplicated neighbors. TEST(AnnIVFFlatTest, CosineZeroNormCenters) { raft::resources handle; auto stream = raft::resource::get_cuda_stream(handle); - constexpr int64_t n_rows = 10000; - constexpr int64_t n_queries = 500; - constexpr int64_t dim = 8; - constexpr uint32_t n_lists = 1024; - constexpr uint32_t k = 16; + constexpr int64_t n_rows = 10000; + constexpr int64_t n_queries = 500; + constexpr int64_t dim = 8; + constexpr uint32_t n_lists = 1024; + constexpr uint32_t n_zero_lists = 1000; + constexpr uint32_t n_probes = 40; + constexpr uint32_t k = 16; auto data = raft::make_device_matrix(handle, n_rows, dim); auto queries = raft::make_device_matrix(handle, n_queries, dim); @@ -76,19 +80,18 @@ TEST(AnnIVFFlatTest, CosineZeroNormCenters) params.add_data_on_build = false; auto idx = build(handle, params, raft::make_const_mdspan(data.view())); - // Zero the centers of every other list, as if k-means had left them empty. + // Zero the centers of lists 1..n_zero_lists, as if k-means had left these clusters empty. Fewer + // lists than n_probes keep a non-zero center, and list 0 is one of them. std::vector centers_h(idx.centers().size()); raft::update_host(centers_h.data(), idx.centers().data_handle(), centers_h.size(), stream); raft::resource::sync_stream(handle); - for (uint32_t l = 0; l < n_lists; l += 2) { - std::fill_n(centers_h.begin() + l * dim, dim, 0.0f); - } + std::fill(centers_h.begin() + dim, centers_h.begin() + (n_zero_lists + 1) * dim, 0.0f); raft::update_device(idx.centers().data_handle(), centers_h.data(), centers_h.size(), stream); extend(handle, raft::make_const_mdspan(data.view()), std::nullopt, &idx); - // Probe exactly as many lists as there are non-zero centers. + // All data rows are in the lists with a non-zero center, and all of those lists are probed. search_params sparams; - sparams.n_probes = n_lists / 2; + sparams.n_probes = n_probes; auto neighbors = raft::make_device_matrix(handle, n_queries, k); auto distances = raft::make_device_matrix(handle, n_queries, k); search(handle, @@ -121,7 +124,6 @@ TEST(AnnIVFFlatTest, CosineZeroNormCenters) naive_distances_h.data(), naive_distances.data_handle(), naive_distances_h.size(), stream); raft::resource::sync_stream(handle); - // Every data row is in a list with a non-zero center, and all of those lists are probed. EXPECT_TRUE(eval_neighbours( naive_neighbors_h, neighbors_h, naive_distances_h, distances_h, n_queries, k, 0.001, 0.99)); } From ebfad2b988b8b83c138e58d25544c7adba91cfc8 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Sun, 11 Oct 2026 03:12:55 +0000 Subject: [PATCH 6/6] [CI ONLY] Compare the zero-norm guard on and off, and run the full IVF-Flat test --- ci/test_cpp.sh | 6 +++++- cpp/src/neighbors/ivf_flat/ivf_flat_search.cuh | 4 ++++ 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/ci/test_cpp.sh b/ci/test_cpp.sh index 5fe18ea172..b26cf537ae 100755 --- a/ci/test_cpp.sh +++ b/ci/test_cpp.sh @@ -70,8 +70,12 @@ repeat_test() { grep -E "^\[ OK \]" "${log}" | sed -E 's/ \([0-9]+ ms\)//' | sort | uniq -c || true } nvidia-smi -q | grep -E "Driver Version|CUDA Version|Product Name" || true +rapids-logger "Zero-norm center regression test WITHOUT the guard (expected to fail)" +CUVS_CI_DISABLE_ZERO_NORM_GUARD=1 "${BIN}" --gtest_filter='AnnIVFFlatTest.CosineZeroNormCenters' 2>&1 | tail -n 20 || true +rapids-logger "Zero-norm center regression test WITH the guard" +"${BIN}" --gtest_filter='AnnIVFFlatTest.CosineZeroNormCenters' --gtest_repeat=20 2>&1 | grep -E "Duplicated|Recall|FAILED|PASSED|OK \]" | sort | uniq -c repeat_test cosine2 '*AnnIVFFlatTestF_int8.AnnIVFFlat/2:*AnnIVFFlatTestF_uint8.AnnIVFFlat/2' 800 -repeat_test int8_uint8_all '*AnnIVFFlatTestF_int8*:*AnnIVFFlatTestF_uint8*' 4 +repeat_test full_ivf_flat '*' 1 rapids-logger "Test script exiting with value: $EXITCODE" exit ${EXITCODE} diff --git a/cpp/src/neighbors/ivf_flat/ivf_flat_search.cuh b/cpp/src/neighbors/ivf_flat/ivf_flat_search.cuh index 4d2d85a8f5..458af0a597 100644 --- a/cpp/src/neighbors/ivf_flat/ivf_flat_search.cuh +++ b/cpp/src/neighbors/ivf_flat/ivf_flat_search.cuh @@ -26,6 +26,7 @@ #include +#include #include namespace cuvs::neighbors::ivf_flat::detail { @@ -165,6 +166,8 @@ void search_impl(raft::resources const& handle, auto n_lists = index.n_lists(); const auto* q_norm_ptr = query_norm_dev.data(); const auto* index_center_norm_ptr = index.center_norms()->data_handle(); + // [CI ONLY] Allow disabling the zero-norm guard to compare against the old behavior. + const bool no_guard = std::getenv("CUVS_CI_DISABLE_ZERO_NORM_GUARD") != nullptr; raft::linalg::map_offset( handle, distance_buffer_dev_view, @@ -175,6 +178,7 @@ void search_impl(raft::resources const& handle, // zero-norm vector is defined as 1 (similarity 0) instead of NaN, which would corrupt the // selection of the lists to probe. const float denom = q_norm_ptr[query] * index_center_norm_ptr[cluster]; + if (no_guard) { return dist / denom; } return denom > 0.0f ? dist / denom : 0.0f; }, raft::make_const_mdspan(distance_buffer_dev_view));