From 43d17b11c8c51c129b439decf5a1dc1a1f426dbd Mon Sep 17 00:00:00 2001 From: achirkin Date: Thu, 19 Jan 2023 13:00:27 +0100 Subject: [PATCH 1/7] Make eval_neighbours fail instead of reporting a warning if the difference against expected recall is not big --- cpp/test/neighbors/ann_ivf_pq.cuh | 8 ++--- cpp/test/neighbors/ann_utils.cuh | 52 +++++++++++++++++-------------- 2 files changed, 32 insertions(+), 28 deletions(-) diff --git a/cpp/test/neighbors/ann_ivf_pq.cuh b/cpp/test/neighbors/ann_ivf_pq.cuh index b5671b74b0..f0f7ee3dc6 100644 --- a/cpp/test/neighbors/ann_ivf_pq.cuh +++ b/cpp/test/neighbors/ann_ivf_pq.cuh @@ -139,8 +139,8 @@ class ivf_pq_test : public ::testing::TestWithParam { protected: void gen_data() { - database.resize(ps.num_db_vecs * ps.dim, stream_); - search_queries.resize(ps.num_queries * ps.dim, stream_); + database.resize(size_t{ps.num_db_vecs} * size_t{ps.dim}, stream_); + search_queries.resize(size_t{ps.num_queries} * size_t{ps.dim}, stream_); raft::random::Rng r(1234ULL); if constexpr (std::is_same{}) { @@ -155,7 +155,7 @@ class ivf_pq_test : public ::testing::TestWithParam { void calc_ref() { - size_t queries_size = ps.num_queries * ps.k; + size_t queries_size = size_t{ps.num_queries} * size_t{ps.k}; rmm::device_uvector distances_naive_dev(queries_size, stream_); rmm::device_uvector indices_naive_dev(queries_size, stream_); naiveBfKnn(distances_naive_dev.data(), @@ -463,7 +463,7 @@ inline auto enum_variety() -> test_cases_t }); ADD_CASE({ x.search_params.lut_dtype = CUDA_R_8U; - x.min_recall = 0.85; + x.min_recall = 0.84; }); ADD_CASE({ diff --git a/cpp/test/neighbors/ann_utils.cuh b/cpp/test/neighbors/ann_utils.cuh index b88b6abd9e..dc2a863589 100644 --- a/cpp/test/neighbors/ann_utils.cuh +++ b/cpp/test/neighbors/ann_utils.cuh @@ -110,28 +110,39 @@ __global__ void naive_distance_kernel(EvalT* dist, IdxT m, IdxT n, IdxT k, - raft::distance::DistanceType type) + raft::distance::DistanceType metric) { - IdxT midx = threadIdx.x + blockIdx.x * blockDim.x; + IdxT midx = IdxT(threadIdx.x) + IdxT(blockIdx.x) * IdxT(blockDim.x); if (midx >= m) return; - for (IdxT nidx = threadIdx.y + blockIdx.y * blockDim.y; nidx < n; - nidx += blockDim.y * gridDim.y) { + IdxT grid_size = IdxT(blockDim.y) * IdxT(gridDim.y); + for (IdxT nidx = threadIdx.y + blockIdx.y * blockDim.y; nidx < n; nidx += grid_size) { EvalT acc = EvalT(0); for (IdxT i = 0; i < k; ++i) { IdxT xidx = i + midx * k; IdxT yidx = i + nidx * k; - EvalT xv = (EvalT)x[xidx]; - EvalT yv = (EvalT)y[yidx]; - if (type == raft::distance::DistanceType::InnerProduct) { - acc += xv * yv; - } else { - EvalT diff = xv - yv; - acc += diff * diff; + auto xv = EvalT(x[xidx]); + auto yv = EvalT(y[yidx]); + switch (metric) { + case raft::distance::DistanceType::InnerProduct: { + acc += xv * yv; + } break; + case raft::distance::DistanceType::L2SqrtExpanded: + case raft::distance::DistanceType::L2SqrtUnexpanded: + case raft::distance::DistanceType::L2Expanded: + case raft::distance::DistanceType::L2Unexpanded: { + auto diff = xv - yv; + acc += diff * diff; + } break; + default: break; } } - if (type == raft::distance::DistanceType::L2SqrtExpanded || - type == raft::distance::DistanceType::L2SqrtUnexpanded) - acc = raft::mySqrt(acc); + switch (metric) { + case raft::distance::DistanceType::L2SqrtExpanded: + case raft::distance::DistanceType::L2SqrtUnexpanded: { + acc = raft::mySqrt(acc); + } break; + default: break; + } dist[midx * n + nidx] = acc; } } @@ -241,16 +252,9 @@ auto eval_neighbours(const std::vector& expected_idx, error_margin < 0 ? "above" : "below", eps); if (actual_recall < min_recall - eps) { - if (actual_recall < min_recall * min_recall - eps) { - RAFT_LOG_ERROR("Recall is much lower than the minimum (%f < %f)", actual_recall, min_recall); - } else { - RAFT_LOG_WARN("Recall is suspiciously too low (%f < %f)", actual_recall, min_recall); - } - if (match_count == 0 || actual_recall < min_recall * std::min(min_recall, 0.5) - eps) { - return testing::AssertionFailure() - << "actual recall (" << actual_recall - << ") is much smaller than the minimum expected recall (" << min_recall << ")."; - } + return testing::AssertionFailure() + << "actual recall (" << actual_recall << ") is lower than the minimum expected recall (" + << min_recall << "); eps = " << eps << ". "; } return testing::AssertionSuccess(); } From d64643a89099ba27d2452351b9610759348d3906 Mon Sep 17 00:00:00 2001 From: achirkin Date: Fri, 20 Jan 2023 10:51:57 +0100 Subject: [PATCH 2/7] Adjust the min_recall for some InnerProduct test cases by 2% --- cpp/test/neighbors/ann_ivf_pq.cuh | 3 +++ 1 file changed, 3 insertions(+) diff --git a/cpp/test/neighbors/ann_ivf_pq.cuh b/cpp/test/neighbors/ann_ivf_pq.cuh index f0f7ee3dc6..eabf971301 100644 --- a/cpp/test/neighbors/ann_ivf_pq.cuh +++ b/cpp/test/neighbors/ann_ivf_pq.cuh @@ -497,6 +497,9 @@ inline auto enum_variety_ip() -> test_cases_t // thus we're forced to used signed 8-bit representation, // thus we have one bit less precision y.min_recall = y.min_recall.value() * 0.95; + } else { + // In other cases it seems to perform just a bit worse than L2 + y.min_recall = y.min_recall.value() * 0.98; } } y.index_params.metric = distance::DistanceType::InnerProduct; From 2350a33c04a7228caaabebd1572e6f3a31e0f49c Mon Sep 17 00:00:00 2001 From: achirkin Date: Fri, 20 Jan 2023 15:01:46 +0100 Subject: [PATCH 3/7] Adjust the min_recall for some InnerProduct test cases by 2% --- cpp/test/neighbors/ann_ivf_pq.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/test/neighbors/ann_ivf_pq.cuh b/cpp/test/neighbors/ann_ivf_pq.cuh index eabf971301..6c47090748 100644 --- a/cpp/test/neighbors/ann_ivf_pq.cuh +++ b/cpp/test/neighbors/ann_ivf_pq.cuh @@ -496,7 +496,7 @@ inline auto enum_variety_ip() -> test_cases_t // InnerProduct score is signed, // thus we're forced to used signed 8-bit representation, // thus we have one bit less precision - y.min_recall = y.min_recall.value() * 0.95; + y.min_recall = y.min_recall.value() * 0.93; } else { // In other cases it seems to perform just a bit worse than L2 y.min_recall = y.min_recall.value() * 0.98; From f7e6dc16f1f0e8896033bf6e8eb3b20458397f1a Mon Sep 17 00:00:00 2001 From: achirkin Date: Fri, 20 Jan 2023 19:08:28 +0100 Subject: [PATCH 4/7] Reduce the min_recall for InnerProduct a little further --- cpp/test/neighbors/ann_ivf_pq.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/test/neighbors/ann_ivf_pq.cuh b/cpp/test/neighbors/ann_ivf_pq.cuh index 6c47090748..8bb2673aa7 100644 --- a/cpp/test/neighbors/ann_ivf_pq.cuh +++ b/cpp/test/neighbors/ann_ivf_pq.cuh @@ -499,7 +499,7 @@ inline auto enum_variety_ip() -> test_cases_t y.min_recall = y.min_recall.value() * 0.93; } else { // In other cases it seems to perform just a bit worse than L2 - y.min_recall = y.min_recall.value() * 0.98; + y.min_recall = y.min_recall.value() * 0.97; } } y.index_params.metric = distance::DistanceType::InnerProduct; From 7b46bcb8ac7c5b6284a72c1456fb5c429ce6b5a9 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Sun, 22 Jan 2023 13:26:49 -0500 Subject: [PATCH 5/7] Fixing bad merge --- cpp/test/neighbors/ann_utils.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/test/neighbors/ann_utils.cuh b/cpp/test/neighbors/ann_utils.cuh index dc2a863589..551ebd767f 100644 --- a/cpp/test/neighbors/ann_utils.cuh +++ b/cpp/test/neighbors/ann_utils.cuh @@ -139,7 +139,7 @@ __global__ void naive_distance_kernel(EvalT* dist, switch (metric) { case raft::distance::DistanceType::L2SqrtExpanded: case raft::distance::DistanceType::L2SqrtUnexpanded: { - acc = raft::mySqrt(acc); + acc = raft::sqrt(acc); } break; default: break; } From 3e4de8d6df7ea6f4079db07be10b9a598eb5e903 Mon Sep 17 00:00:00 2001 From: achirkin Date: Mon, 23 Jan 2023 07:22:06 +0100 Subject: [PATCH 6/7] Further adjust InnerProduct thresholds --- cpp/test/neighbors/ann_ivf_pq.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/test/neighbors/ann_ivf_pq.cuh b/cpp/test/neighbors/ann_ivf_pq.cuh index 8bb2673aa7..c67feae522 100644 --- a/cpp/test/neighbors/ann_ivf_pq.cuh +++ b/cpp/test/neighbors/ann_ivf_pq.cuh @@ -498,8 +498,8 @@ inline auto enum_variety_ip() -> test_cases_t // thus we have one bit less precision y.min_recall = y.min_recall.value() * 0.93; } else { - // In other cases it seems to perform just a bit worse than L2 - y.min_recall = y.min_recall.value() * 0.97; + // In other cases it seems to perform a little bit better, still worse than L2 + y.min_recall = y.min_recall.value() * 0.95; } } y.index_params.metric = distance::DistanceType::InnerProduct; From d317b6ed9323cadc7d26b61555d4c3adefbe0688 Mon Sep 17 00:00:00 2001 From: achirkin Date: Mon, 23 Jan 2023 10:24:43 +0100 Subject: [PATCH 7/7] Adjusted InnerProduct threshold further after more testing on A100 --- cpp/test/neighbors/ann_ivf_pq.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/test/neighbors/ann_ivf_pq.cuh b/cpp/test/neighbors/ann_ivf_pq.cuh index c67feae522..719f429f13 100644 --- a/cpp/test/neighbors/ann_ivf_pq.cuh +++ b/cpp/test/neighbors/ann_ivf_pq.cuh @@ -496,10 +496,10 @@ inline auto enum_variety_ip() -> test_cases_t // InnerProduct score is signed, // thus we're forced to used signed 8-bit representation, // thus we have one bit less precision - y.min_recall = y.min_recall.value() * 0.93; + y.min_recall = y.min_recall.value() * 0.90; } else { // In other cases it seems to perform a little bit better, still worse than L2 - y.min_recall = y.min_recall.value() * 0.95; + y.min_recall = y.min_recall.value() * 0.94; } } y.index_params.metric = distance::DistanceType::InnerProduct;