From bd694e9d01e4459118d130545ad962b099c904d6 Mon Sep 17 00:00:00 2001 From: Yan Zaretskiy Date: Thu, 19 Feb 2026 17:57:54 -0800 Subject: [PATCH 1/3] L1 distance support for iterative search CAGRA build --- cpp/CMakeLists.txt | 12 +++++ .../detail/cagra/compute_distance-ext.cuh | 48 ++++++++++++++++++- .../detail/cagra/compute_distance.cu | 14 +++++- .../cagra/compute_distance_00_generate.py | 12 ++--- .../cagra/compute_distance_standard-impl.cuh | 21 ++++---- ...ance_standard_L1_float_uint32_dim128_t8.cu | 22 +++++++++ ...nce_standard_L1_float_uint32_dim256_t16.cu | 22 +++++++++ ...nce_standard_L1_float_uint32_dim512_t32.cu | 22 +++++++++ ...tance_standard_L1_half_uint32_dim128_t8.cu | 22 +++++++++ ...ance_standard_L1_half_uint32_dim256_t16.cu | 22 +++++++++ ...ance_standard_L1_half_uint32_dim512_t32.cu | 22 +++++++++ ...tance_standard_L1_int8_uint32_dim128_t8.cu | 22 +++++++++ ...ance_standard_L1_int8_uint32_dim256_t16.cu | 22 +++++++++ ...ance_standard_L1_int8_uint32_dim512_t32.cu | 22 +++++++++ ...ance_standard_L1_uint8_uint32_dim128_t8.cu | 22 +++++++++ ...nce_standard_L1_uint8_uint32_dim256_t16.cu | 22 +++++++++ ...nce_standard_L1_uint8_uint32_dim512_t32.cu | 22 +++++++++ cpp/src/neighbors/detail/cagra/graph_core.cuh | 13 ++++- cpp/src/neighbors/ivf_common.cuh | 12 +++++ cpp/tests/neighbors/ann_cagra.cuh | 37 +++++++++++--- cpp/tests/neighbors/naive_knn.cuh | 6 ++- 21 files changed, 413 insertions(+), 26 deletions(-) create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim128_t8.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim256_t16.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim512_t32.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim128_t8.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim256_t16.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim512_t32.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim128_t8.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim256_t16.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim512_t32.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim128_t8.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim256_t16.cu create mode 100644 cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim512_t32.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index a75890737e..501dfb87c8 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -249,6 +249,18 @@ if(NOT BUILD_CPU_ONLY) src/neighbors/detail/cagra/compute_distance_standard_InnerProduct_uint8_uint32_dim128_t8.cu src/neighbors/detail/cagra/compute_distance_standard_InnerProduct_uint8_uint32_dim256_t16.cu src/neighbors/detail/cagra/compute_distance_standard_InnerProduct_uint8_uint32_dim512_t32.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim128_t8.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim256_t16.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim512_t32.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim128_t8.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim256_t16.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim512_t32.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim128_t8.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim256_t16.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim512_t32.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim128_t8.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim256_t16.cu + src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim512_t32.cu src/neighbors/detail/cagra/compute_distance_standard_L2Expanded_float_uint32_dim128_t8.cu src/neighbors/detail/cagra/compute_distance_standard_L2Expanded_float_uint32_dim256_t16.cu src/neighbors/detail/cagra/compute_distance_standard_L2Expanded_float_uint32_dim512_t32.cu diff --git a/cpp/src/neighbors/detail/cagra/compute_distance-ext.cuh b/cpp/src/neighbors/detail/cagra/compute_distance-ext.cuh index f0eababe36..7ce684ac73 100644 --- a/cpp/src/neighbors/detail/cagra/compute_distance-ext.cuh +++ b/cpp/src/neighbors/detail/cagra/compute_distance-ext.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -39,6 +39,7 @@ extern template struct standard_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec; +extern template struct standard_descriptor_spec; extern template struct vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, @@ -541,61 +575,73 @@ using descriptor_instances = instance_selector< standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, diff --git a/cpp/src/neighbors/detail/cagra/compute_distance.cu b/cpp/src/neighbors/detail/cagra/compute_distance.cu index a0b7209814..76a299921a 100644 --- a/cpp/src/neighbors/detail/cagra/compute_distance.cu +++ b/cpp/src/neighbors/detail/cagra/compute_distance.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -22,61 +22,73 @@ template struct instance_selector< standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, standard_descriptor_spec, + standard_descriptor_spec, vpq_descriptor_spec, vpq_descriptor_spec, standard_descriptor_spec, diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_00_generate.py b/cpp/src/neighbors/detail/cagra/compute_distance_00_generate.py index fde2081c12..3892e4dd42 100644 --- a/cpp/src/neighbors/detail/cagra/compute_distance_00_generate.py +++ b/cpp/src/neighbors/detail/cagra/compute_distance_00_generate.py @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 import datetime import os @@ -19,14 +19,14 @@ * */ -{{includes}} +{includes} -namespace cuvs::neighbors::cagra::detail {{{{ +namespace cuvs::neighbors::cagra::detail {{ using namespace cuvs::distance; -{{content}} +{content} -}}}} // namespace cuvs::neighbors::cagra::detail +}} // namespace cuvs::neighbors::cagra::detail """ mxdim_team = [(128, 8), (256, 16), (512, 32)] @@ -65,7 +65,7 @@ for type_path, (data_t, idx_t, distance_t) in search_types.items(): for mxdim, team in mxdim_team: # CAGRA - for metric in ["L2Expanded", "InnerProduct", "CosineExpanded"]: + for metric in ["L2Expanded", "InnerProduct", "CosineExpanded", "L1"]: path = f"compute_distance_standard_{metric}_{type_path}_dim{mxdim}_t{team}.cu" includes = '#include "compute_distance_standard-impl.cuh"' params = f"{metric_prefix}{metric}, {team}, {mxdim}, {data_t}, {idx_t}, {distance_t}" diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard-impl.cuh b/cpp/src/neighbors/detail/cagra/compute_distance_standard-impl.cuh index ecb09f516c..05adce20e9 100644 --- a/cpp/src/neighbors/detail/cagra/compute_distance_standard-impl.cuh +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard-impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -15,32 +15,37 @@ namespace cuvs::neighbors::cagra::detail { namespace { template + requires(Metric == cuvs::distance::DistanceType::L2Expanded) RAFT_DEVICE_INLINE_FUNCTION constexpr auto dist_op(DATA_T a, DATA_T b) - -> std::enable_if_t { DISTANCE_T diff = a - b; return diff * diff; } template + requires(Metric == cuvs::distance::DistanceType::InnerProduct || + Metric == cuvs::distance::DistanceType::CosineExpanded) RAFT_DEVICE_INLINE_FUNCTION constexpr auto dist_op(DATA_T a, DATA_T b) - -> std::enable_if_t { return -static_cast(a) * static_cast(b); } template + requires(Metric == cuvs::distance::DistanceType::BitwiseHamming && std::is_integral_v) RAFT_DEVICE_INLINE_FUNCTION constexpr auto dist_op(DATA_T a, DATA_T b) - -> std::enable_if_t, - DISTANCE_T> { // mask the result of xor for the integer promotion const auto v = (a ^ b) & 0xffu; return __popc(v); } + +template + requires(Metric == cuvs::distance::DistanceType::L1) +RAFT_DEVICE_INLINE_FUNCTION constexpr auto dist_op(DATA_T a, DATA_T b) +{ + DISTANCE_T diff = a - b; + return raft::abs(diff); +} } // namespace template python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim256_t16.cu new file mode 100644 index 0000000000..96a6c8e2bd --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim256_t16.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim512_t32.cu new file mode 100644 index 0000000000..6efb65e9b3 --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_float_uint32_dim512_t32.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim128_t8.cu new file mode 100644 index 0000000000..0f099a0ed4 --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim128_t8.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim256_t16.cu new file mode 100644 index 0000000000..3fa1503e97 --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim256_t16.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim512_t32.cu new file mode 100644 index 0000000000..605e11859d --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_half_uint32_dim512_t32.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim128_t8.cu new file mode 100644 index 0000000000..c75883a3a9 --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim128_t8.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim256_t16.cu new file mode 100644 index 0000000000..a0eb415c39 --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim256_t16.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim512_t32.cu new file mode 100644 index 0000000000..92672cb059 --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_int8_uint32_dim512_t32.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim128_t8.cu new file mode 100644 index 0000000000..02f8ace308 --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim128_t8.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim256_t16.cu new file mode 100644 index 0000000000..5315729602 --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim256_t16.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim512_t32.cu new file mode 100644 index 0000000000..be6703e9f4 --- /dev/null +++ b/cpp/src/neighbors/detail/cagra/compute_distance_standard_L1_uint8_uint32_dim512_t32.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +/* + * NOTE: this file is generated by compute_distance_00_generate.py + * + * Make changes there and run in this directory: + * + * > python compute_distance_00_generate.py + * + */ + +#include "compute_distance_standard-impl.cuh" + +namespace cuvs::neighbors::cagra::detail { + +using namespace cuvs::distance; +template struct standard_descriptor_spec; + +} // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/graph_core.cuh b/cpp/src/neighbors/detail/cagra/graph_core.cuh index 41efa1686f..2a4857bd83 100644 --- a/cpp/src/neighbors/detail/cagra/graph_core.cuh +++ b/cpp/src/neighbors/detail/cagra/graph_core.cuh @@ -105,6 +105,14 @@ __global__ void kern_sort(const DATA_T* const dataset, // [dataset_chunk_size, dataset[d + static_cast(dataset_dim) * dstNode]); dist += diff * diff; } + } else if (metric == cuvs::distance::DistanceType::L1) { + for (int d = lane_id; d < dataset_dim; d += raft::WarpSize) { + float diff = cuvs::spatial::knn::detail::utils::mapping{}( + dataset[d + static_cast(dataset_dim) * srcNode]) - + cuvs::spatial::knn::detail::utils::mapping{}( + dataset[d + static_cast(dataset_dim) * dstNode]); + dist += raft::abs(diff); + } } else if (metric == cuvs::distance::DistanceType::BitwiseHamming) { if constexpr (std::is_integral_v) { for (int d = lane_id; d < dataset_dim; d += raft::WarpSize) { @@ -507,8 +515,9 @@ void sort_knn_graph( metric == cuvs::distance::DistanceType::InnerProduct || metric == cuvs::distance::DistanceType::CosineExpanded || metric == cuvs::distance::DistanceType::L2Expanded || - metric == cuvs::distance::DistanceType::BitwiseHamming, - "Unsupported metric. Only InnerProduct, CosineExpanded, L2Expanded and BitwiseHamming are " + metric == cuvs::distance::DistanceType::BitwiseHamming || + metric == cuvs::distance::DistanceType::L1, + "Unsupported metric. Only InnerProduct, CosineExpanded, L2Expanded, BitwiseHamming and L1 are " "supported"); const uint64_t dataset_size = dataset.extent(0); const uint64_t dataset_dim = dataset.extent(1); diff --git a/cpp/src/neighbors/ivf_common.cuh b/cpp/src/neighbors/ivf_common.cuh index ad3dc86d0d..f91f23a3b9 100644 --- a/cpp/src/neighbors/ivf_common.cuh +++ b/cpp/src/neighbors/ivf_common.cuh @@ -228,6 +228,18 @@ void postprocess_distances(ScoreOutT* out, // [n_queries, topk] } } break; case distance::DistanceType::BitwiseHamming: break; + case distance::DistanceType::L1: { + if (scaling_factor != 1.0) { + raft::linalg::unaryOp(out, + in, + len, + raft::compose_op(raft::mul_const_op{scaling_factor}, + raft::cast_op{}), + stream); + } else if (needs_cast || needs_copy) { + raft::linalg::unaryOp(out, in, len, raft::cast_op{}, stream); + } + } break; default: RAFT_FAIL("Unexpected metric."); } } diff --git a/cpp/tests/neighbors/ann_cagra.cuh b/cpp/tests/neighbors/ann_cagra.cuh index beb379e44d..ea82e86c0c 100644 --- a/cpp/tests/neighbors/ann_cagra.cuh +++ b/cpp/tests/neighbors/ann_cagra.cuh @@ -286,6 +286,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, const AnnCagraInputs& p) case cuvs::distance::DistanceType::L2Expanded: return "L2"; case cuvs::distance::DistanceType::BitwiseHamming: return "BitwiseHamming"; case cuvs::distance::DistanceType::CosineExpanded: return "Cosine"; + case cuvs::distance::DistanceType::L1: return "L1"; default: break; } return "Unknown"; @@ -342,6 +343,9 @@ class AnnCagraTest : public ::testing::TestWithParam { if (ps.metric == cuvs::distance::DistanceType::BitwiseHamming && (ps.k * ps.dim * 8 / 5 /*(=magic number)*/ < ps.n_rows)) GTEST_SKIP(); + if (ps.metric == cuvs::distance::DistanceType::L1 && + ps.build_algo != graph_build_algo::ITERATIVE_CAGRA_SEARCH) + GTEST_SKIP(); if (ps.metric == cuvs::distance::DistanceType::CosineExpanded) { if (ps.compression.has_value()) { GTEST_SKIP(); } if (ps.build_algo == graph_build_algo::ITERATIVE_CAGRA_SEARCH || ps.dim == 1) { @@ -532,6 +536,9 @@ class AnnCagraAddNodesTest : public ::testing::TestWithParam { protected: void testCagra() { + if (ps.metric == cuvs::distance::DistanceType::L1 && + ps.build_algo != graph_build_algo::ITERATIVE_CAGRA_SEARCH) + GTEST_SKIP(); if (ps.metric == cuvs::distance::DistanceType::CosineExpanded) { if (ps.compression.has_value()) { GTEST_SKIP(); } if (ps.build_algo == graph_build_algo::ITERATIVE_CAGRA_SEARCH || ps.dim == 1) { @@ -745,6 +752,9 @@ class AnnCagraFilterTest : public ::testing::TestWithParam { protected: void testCagra() { + if (ps.metric == cuvs::distance::DistanceType::L1 && + ps.build_algo != graph_build_algo::ITERATIVE_CAGRA_SEARCH) + GTEST_SKIP(); if (ps.metric == cuvs::distance::DistanceType::CosineExpanded) { if (ps.compression.has_value()) { GTEST_SKIP(); } if (ps.build_algo == graph_build_algo::ITERATIVE_CAGRA_SEARCH || ps.dim == 1) { @@ -964,6 +974,9 @@ class AnnCagraIndexFilteredMergeTest : public ::testing::TestWithParam void testCagra() { + if (ps.metric == cuvs::distance::DistanceType::L1 && + ps.build_algo != graph_build_algo::ITERATIVE_CAGRA_SEARCH) + GTEST_SKIP(); if (ps.metric == cuvs::distance::DistanceType::CosineExpanded) { if (ps.build_algo == graph_build_algo::ITERATIVE_CAGRA_SEARCH || ps.dim == 1) { GTEST_SKIP(); @@ -1205,6 +1218,9 @@ class AnnCagraIndexMergeTest : public ::testing::TestWithParam { template void testCagra() { + if (ps.metric == cuvs::distance::DistanceType::L1 && + ps.build_algo != graph_build_algo::ITERATIVE_CAGRA_SEARCH) + GTEST_SKIP(); if (ps.metric == cuvs::distance::DistanceType::CosineExpanded) { if (ps.build_algo == graph_build_algo::ITERATIVE_CAGRA_SEARCH || ps.dim == 1) { GTEST_SKIP(); @@ -1419,7 +1435,8 @@ inline std::vector generate_inputs() {cuvs::distance::DistanceType::L2Expanded, cuvs::distance::DistanceType::InnerProduct, cuvs::distance::DistanceType::BitwiseHamming, - cuvs::distance::DistanceType::CosineExpanded}, + cuvs::distance::DistanceType::CosineExpanded, + cuvs::distance::DistanceType::L1}, {false}, {true}, {true, false}, @@ -1443,7 +1460,8 @@ inline std::vector generate_inputs() {cuvs::distance::DistanceType::L2Expanded, cuvs::distance::DistanceType::InnerProduct, cuvs::distance::DistanceType::BitwiseHamming, - cuvs::distance::DistanceType::CosineExpanded}, + cuvs::distance::DistanceType::CosineExpanded, + cuvs::distance::DistanceType::L1}, {false}, {true}, {false}, @@ -1468,7 +1486,8 @@ inline std::vector generate_inputs() {1}, {cuvs::distance::DistanceType::InnerProduct, cuvs::distance::DistanceType::BitwiseHamming, - cuvs::distance::DistanceType::CosineExpanded}, + cuvs::distance::DistanceType::CosineExpanded, + cuvs::distance::DistanceType::L1}, {false}, {true}, {false}, @@ -1519,7 +1538,8 @@ inline std::vector generate_inputs() {1}, {cuvs::distance::DistanceType::L2Expanded, cuvs::distance::DistanceType::InnerProduct, - cuvs::distance::DistanceType::BitwiseHamming}, + cuvs::distance::DistanceType::BitwiseHamming, + cuvs::distance::DistanceType::L1}, {false}, {true}, {false}, @@ -1548,7 +1568,8 @@ inline std::vector generate_inputs() {cuvs::distance::DistanceType::L2Expanded, cuvs::distance::DistanceType::InnerProduct, cuvs::distance::DistanceType::BitwiseHamming, - cuvs::distance::DistanceType::CosineExpanded}, + cuvs::distance::DistanceType::CosineExpanded, + cuvs::distance::DistanceType::L1}, {false}, {false}, {false}, @@ -1575,7 +1596,8 @@ inline std::vector generate_inputs() {cuvs::distance::DistanceType::L2Expanded, cuvs::distance::DistanceType::InnerProduct, cuvs::distance::DistanceType::BitwiseHamming, - cuvs::distance::DistanceType::CosineExpanded}, + cuvs::distance::DistanceType::CosineExpanded, + cuvs::distance::DistanceType::L1}, {false}, {false}, {false}, @@ -1717,7 +1739,8 @@ inline std::vector generate_addnode_inputs() {1}, {cuvs::distance::DistanceType::L2Expanded, cuvs::distance::DistanceType::InnerProduct, - cuvs::distance::DistanceType::BitwiseHamming}, + cuvs::distance::DistanceType::BitwiseHamming, + cuvs::distance::DistanceType::L1}, {false}, {true}, {true}, diff --git a/cpp/tests/neighbors/naive_knn.cuh b/cpp/tests/neighbors/naive_knn.cuh index 0b549cfdb1..2ed461360c 100644 --- a/cpp/tests/neighbors/naive_knn.cuh +++ b/cpp/tests/neighbors/naive_knn.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -59,6 +59,10 @@ RAFT_KERNEL naive_distance_kernel(EvalT* dist, acc += __popc(static_cast(xv ^ yv) & 0xff); } } break; + case cuvs::distance::DistanceType::L1: { + auto diff = static_cast(xv) - static_cast(yv); + acc += raft::abs(diff); + } break; default: break; } } From 135c05b3a439a78011c4e313dd094ca751b466cc Mon Sep 17 00:00:00 2001 From: Yan Zaretskiy Date: Thu, 5 Mar 2026 19:53:40 +0000 Subject: [PATCH 2/3] Fix the postprocess_distances L1 branch to match the recent main branch refactor --- cpp/src/neighbors/ivf_common.cuh | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/cpp/src/neighbors/ivf_common.cuh b/cpp/src/neighbors/ivf_common.cuh index f1b88ab132..f441e6326e 100644 --- a/cpp/src/neighbors/ivf_common.cuh +++ b/cpp/src/neighbors/ivf_common.cuh @@ -237,14 +237,15 @@ void postprocess_distances(const raft::resources& res, case distance::DistanceType::BitwiseHamming: break; case distance::DistanceType::L1: { if (scaling_factor != 1.0) { - raft::linalg::unaryOp(out, - in, - len, - raft::compose_op(raft::mul_const_op{scaling_factor}, - raft::cast_op{}), - stream); + raft::linalg::map( + res, + out_view, + raft::compose_op(raft::mul_const_op{scaling_factor}, + raft::cast_op{}), + raft::make_const_mdspan(in_view)); } else if (needs_cast || needs_copy) { - raft::linalg::unaryOp(out, in, len, raft::cast_op{}, stream); + raft::linalg::map( + res, out_view, raft::cast_op{}, raft::make_const_mdspan(in_view)); } } break; default: RAFT_FAIL("Unexpected metric."); From f4fa07847986399699f5ecb5b5dacfa8dc294adf Mon Sep 17 00:00:00 2001 From: Yan Zaretskiy Date: Thu, 5 Mar 2026 20:15:39 +0000 Subject: [PATCH 3/3] Fix pre-commit formatting check failure --- cpp/src/neighbors/ivf_common.cuh | 11 +++++------ 1 file changed, 5 insertions(+), 6 deletions(-) diff --git a/cpp/src/neighbors/ivf_common.cuh b/cpp/src/neighbors/ivf_common.cuh index f441e6326e..80aac970dd 100644 --- a/cpp/src/neighbors/ivf_common.cuh +++ b/cpp/src/neighbors/ivf_common.cuh @@ -237,12 +237,11 @@ void postprocess_distances(const raft::resources& res, case distance::DistanceType::BitwiseHamming: break; case distance::DistanceType::L1: { if (scaling_factor != 1.0) { - raft::linalg::map( - res, - out_view, - raft::compose_op(raft::mul_const_op{scaling_factor}, - raft::cast_op{}), - raft::make_const_mdspan(in_view)); + raft::linalg::map(res, + out_view, + raft::compose_op(raft::mul_const_op{scaling_factor}, + raft::cast_op{}), + raft::make_const_mdspan(in_view)); } else if (needs_cast || needs_copy) { raft::linalg::map( res, out_view, raft::cast_op{}, raft::make_const_mdspan(in_view));