From dac4f92bd7910f5f277458f27516bf6960c6267b Mon Sep 17 00:00:00 2001 From: Divye Gala Date: Fri, 23 Oct 2026 00:46:13 -0700 Subject: [PATCH 1/2] Share CUB histogram launchers --- cpp/CMakeLists.txt | 20 +++++ cpp/src/cluster/detail/kmeans_common.cuh | 75 ++++++++++------- cpp/src/core/count_labels.cu.in | 49 +++++++++++ cpp/src/core/count_labels_matrix.json | 26 ++++++ cpp/src/core/cub_histogram.hpp | 36 ++++++++ cpp/src/core/histogram_even.cu.in | 34 ++++++++ cpp/src/core/histogram_even_matrix.json | 82 +++++++++++++++++++ .../neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu | 73 +++++++++-------- cpp/src/neighbors/scann/detail/scann_avq.cuh | 37 +++++---- cpp/src/stats/detail/silhouette_score.cuh | 37 +++++---- 10 files changed, 366 insertions(+), 103 deletions(-) create mode 100644 cpp/src/core/count_labels.cu.in create mode 100644 cpp/src/core/count_labels_matrix.json create mode 100644 cpp/src/core/cub_histogram.hpp create mode 100644 cpp/src/core/histogram_even.cu.in create mode 100644 cpp/src/core/histogram_even_matrix.json diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index ba8a7c7353..6baaf073d6 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -1431,8 +1431,28 @@ if(NOT BUILD_CPU_ONLY) "${CMAKE_CURRENT_BINARY_DIR}/src/selection/select_k_inst_data_@data_abbrev@_index_@index_abbrev@.cu" ) + # Keep every primitive function in its own CUDA translation unit. + foreach(primitive IN ITEMS src/core/count_labels src/core/histogram_even) + set(primitive_instantiations) + generate_string_matrix( + primitive_instantiations + MATRIX_JSON_FILE + "${CMAKE_CURRENT_SOURCE_DIR}/${primitive}_matrix.json" + ITEM_FORMAT + "template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@@semicolon@" + GLUE + "\n" + ) + configure_file( + "${CMAKE_CURRENT_SOURCE_DIR}/${primitive}.cu.in" + "${CMAKE_CURRENT_BINARY_DIR}/${primitive}.cu" @ONLY + ) + list(APPEND cuvs_primitive_sources "${CMAKE_CURRENT_BINARY_DIR}/${primitive}.cu") + endforeach() + add_library( cuvs_objs OBJECT + ${cuvs_primitive_sources} src/core/cub_ptx_query.cu src/cluster/detail/minClusterDistanceCompute.cu src/cluster/agglomerative.cu diff --git a/cpp/src/cluster/detail/kmeans_common.cuh b/cpp/src/cluster/detail/kmeans_common.cuh index 78edf8759e..100518b20e 100644 --- a/cpp/src/cluster/detail/kmeans_common.cuh +++ b/cpp/src/cluster/detail/kmeans_common.cuh @@ -4,6 +4,7 @@ */ #pragma once +#include "../../core/cub_histogram.hpp" #include "../../distance/distance.cuh" #include "../../distance/top_1_nn.cuh" #include @@ -143,37 +144,49 @@ void countLabels(raft::resources const& handle, IndexT n_clusters, rmm::device_uvector& workspace) { - cudaStream_t stream = raft::resource::get_cuda_stream(handle).get(); - - // CUB::DeviceHistogram requires a signed index type - typedef typename std::make_signed_t CubIndexT; - - CubIndexT num_levels = n_clusters + 1; - CubIndexT lower_level = 0; - CubIndexT upper_level = n_clusters; - - size_t temp_storage_bytes = 0; - RAFT_CUDA_TRY(cub::DeviceHistogram::HistogramEven(nullptr, - temp_storage_bytes, - labels, - count, - num_levels, - lower_level, - upper_level, - static_cast(n_samples), - stream)); - - workspace.resize(temp_storage_bytes, stream); - - RAFT_CUDA_TRY(cub::DeviceHistogram::HistogramEven(workspace.data(), - temp_storage_bytes, - labels, - count, - num_levels, - lower_level, - upper_level, - static_cast(n_samples), - stream)); + if constexpr (std::is_same_v && + (std::is_same_v || + std::is_same_v)) { + using CubIndexT = std::make_signed_t; + cuvs::detail::count_labels(labels, + count, + static_cast(n_samples), + static_cast(n_clusters), + workspace, + raft::resource::get_cuda_stream(handle).get()); + } else { + cudaStream_t stream = raft::resource::get_cuda_stream(handle).get(); + + // CUB::DeviceHistogram requires a signed index type + typedef typename std::make_signed_t CubIndexT; + + CubIndexT num_levels = n_clusters + 1; + CubIndexT lower_level = 0; + CubIndexT upper_level = n_clusters; + + size_t temp_storage_bytes = 0; + RAFT_CUDA_TRY(cuvs::detail::histogram_even(nullptr, + temp_storage_bytes, + labels, + count, + num_levels, + lower_level, + upper_level, + static_cast(n_samples), + stream)); + + workspace.resize(temp_storage_bytes, stream); + + RAFT_CUDA_TRY(cuvs::detail::histogram_even(workspace.data(), + temp_storage_bytes, + labels, + count, + num_levels, + lower_level, + upper_level, + static_cast(n_samples), + stream)); + } } /** diff --git a/cpp/src/core/count_labels.cu.in b/cpp/src/core/count_labels.cu.in new file mode 100644 index 0000000000..532bc84761 --- /dev/null +++ b/cpp/src/core/count_labels.cu.in @@ -0,0 +1,49 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include +#include +#include + +namespace cuvs::detail { + +template +void count_labels(const uint32_t* labels, + CounterT* counts, + IndexT n_samples, + IndexT n_clusters, + rmm::device_uvector& workspace, + cudaStream_t stream) +{ + IndexT num_levels = n_clusters + 1; + IndexT lower_level = 0; + IndexT upper_level = n_clusters; + size_t temp_storage_bytes = 0; + RAFT_CUDA_TRY(cuvs::detail::histogram_even(nullptr, + temp_storage_bytes, + labels, + counts, + num_levels, + lower_level, + upper_level, + n_samples, + stream)); + workspace.resize(temp_storage_bytes, stream); + RAFT_CUDA_TRY(cuvs::detail::histogram_even(workspace.data(), + temp_storage_bytes, + labels, + counts, + num_levels, + lower_level, + upper_level, + n_samples, + stream)); +} + +@primitive_instantiations@ + +} // namespace cuvs::detail diff --git a/cpp/src/core/count_labels_matrix.json b/cpp/src/core/count_labels_matrix.json new file mode 100644 index 0000000000..78fdd81656 --- /dev/null +++ b/cpp/src/core/count_labels_matrix.json @@ -0,0 +1,26 @@ +[ + { + "return_type": "void", + "qualified_function": "count_labels", + "function_qualifiers": "", + "namespace_context": "cuvs::detail", + "_variants": [ + { + "template_arguments": "", + "parameter_types": "const uint32_t*, uint32_t*, int, int, rmm::device_uvector&, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "const uint32_t*, uint32_t*, int64_t, int64_t, rmm::device_uvector&, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "const uint32_t*, unsigned long long*, int, int, rmm::device_uvector&, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "const uint32_t*, unsigned long long*, int64_t, int64_t, rmm::device_uvector&, cudaStream_t" + } + ] + } +] diff --git a/cpp/src/core/cub_histogram.hpp b/cpp/src/core/cub_histogram.hpp new file mode 100644 index 0000000000..12ad63fc67 --- /dev/null +++ b/cpp/src/core/cub_histogram.hpp @@ -0,0 +1,36 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include +#include + +namespace cuvs::detail { + +/** CUB label histograms shared by balanced k-means instantiations. */ +template +void count_labels(const uint32_t* labels, + CounterT* counts, + IndexT n_samples, + IndexT n_clusters, + rmm::device_uvector& workspace, + cudaStream_t stream); + +} // namespace cuvs::detail + +namespace cuvs::detail { +template +cudaError_t histogram_even(void* storage, + size_t& bytes, + SampleIt samples, + CounterT* counts, + int levels, + LevelT lower, + LevelT upper, + NumItems count, + cudaStream_t stream); +} // namespace cuvs::detail diff --git a/cpp/src/core/histogram_even.cu.in b/cpp/src/core/histogram_even.cu.in new file mode 100644 index 0000000000..698a501816 --- /dev/null +++ b/cpp/src/core/histogram_even.cu.in @@ -0,0 +1,34 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include +#include +#include +#include +#include +#include + +namespace cuvs::detail { + +template +cudaError_t histogram_even(void* storage, + size_t& bytes, + SampleIt samples, + CounterT* counts, + int levels, + LevelT lower, + LevelT upper, + NumItems count, + cudaStream_t stream) +{ + return cub::DeviceHistogram::HistogramEven( + storage, bytes, samples, counts, levels, lower, upper, count, stream); +} + +@primitive_instantiations@ + +} // namespace cuvs::detail diff --git a/cpp/src/core/histogram_even_matrix.json b/cpp/src/core/histogram_even_matrix.json new file mode 100644 index 0000000000..5dcb79543f --- /dev/null +++ b/cpp/src/core/histogram_even_matrix.json @@ -0,0 +1,82 @@ +[ + { + "return_type": "cudaError_t", + "qualified_function": "cuvs::detail::histogram_even", + "function_qualifiers": "", + "namespace_context": "cuvs::detail", + "_variants": [ + { + "template_arguments": "< ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, double, int, int>", + "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, double*, int, int, int, int, cudaStream_t" + }, + { + "template_arguments": "< ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, float, int, int>", + "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, float*, int, int, int, int, cudaStream_t" + }, + { + "template_arguments": "< ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, double, long, long>", + "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, double*, int, long, long, long, cudaStream_t" + }, + { + "template_arguments": "< ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, float, long, long>", + "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, float*, int, long, long, long, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, int const*, double*, int, int, int, int, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, int const*, float*, int, int, int, int, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, int const*, int*, int, int, int, int, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, int*, double*, int, int, int, int, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, int*, float*, int, int, int, int, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, int*, int*, int, int, int, int, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, long*, double*, int, long, long, long, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, long*, float*, int, long, long, long, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned int*, int, int, int, long, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned long long*, int, int, int, unsigned long, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned int*, int, int, int, int, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned int*, int, long, long, long, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned long long*, int, int, int, int, cudaStream_t" + }, + { + "template_arguments": "", + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned long long*, int, long, long, long, cudaStream_t" + } + ] + } +] diff --git a/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu b/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu index 9ce10fd57e..4ae5a740c0 100644 --- a/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu +++ b/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu @@ -7,6 +7,7 @@ // Created by Stardust on 2/23/25. // +#include "../../../core/cub_histogram.hpp" #include "../../../core/omp_wrapper.hpp" #include "../../../util/serialize_validation.hpp" #include "../../detail/ann_utils.cuh" @@ -428,27 +429,27 @@ void IVFGPU::construct_on_gpu(const float* device_data, // Determine temporary device storage requirements size_t temp_storage_bytes = 0; - RAFT_CUDA_TRY(cub::DeviceHistogram::HistogramEven(nullptr, - temp_storage_bytes, - device_cluster_ids, - d_histogram.data(), - num_levels, - lower_level, - upper_level, - num_vectors, - stream_.get())); + RAFT_CUDA_TRY(cuvs::detail::histogram_even(nullptr, + temp_storage_bytes, + device_cluster_ids, + d_histogram.data(), + num_levels, + lower_level, + upper_level, + num_vectors, + stream_.get())); { rmm::device_buffer d_temp_storage(temp_storage_bytes, stream_); - RAFT_CUDA_TRY(cub::DeviceHistogram::HistogramEven(d_temp_storage.data(), - temp_storage_bytes, - device_cluster_ids, - d_histogram.data(), - num_levels, - lower_level, - upper_level, - num_vectors, - stream_.get())); + RAFT_CUDA_TRY(cuvs::detail::histogram_even(d_temp_storage.data(), + temp_storage_bytes, + device_cluster_ids, + d_histogram.data(), + num_levels, + lower_level, + upper_level, + num_vectors, + stream_.get())); } // ------------------------- @@ -585,28 +586,28 @@ void IVFGPU::construct_on_gpu_streaming(const float* host_data, void* d_temp_storage = nullptr; size_t temp_storage_bytes = 0; - RAFT_CUDA_TRY(cub::DeviceHistogram::HistogramEven(d_temp_storage, - temp_storage_bytes, - device_cluster_ids, - d_histogram.data_handle(), - num_levels, - lower_level, - upper_level, - num_vectors, - stream_.get())); + RAFT_CUDA_TRY(cuvs::detail::histogram_even(d_temp_storage, + temp_storage_bytes, + device_cluster_ids, + d_histogram.data_handle(), + num_levels, + lower_level, + upper_level, + num_vectors, + stream_.get())); auto d_temp_storage_vec = raft::make_device_vector(handle_, temp_storage_bytes); d_temp_storage = d_temp_storage_vec.data_handle(); - RAFT_CUDA_TRY(cub::DeviceHistogram::HistogramEven(d_temp_storage, - temp_storage_bytes, - device_cluster_ids, - d_histogram.data_handle(), - num_levels, - lower_level, - upper_level, - num_vectors, - stream_.get())); + RAFT_CUDA_TRY(cuvs::detail::histogram_even(d_temp_storage, + temp_storage_bytes, + device_cluster_ids, + d_histogram.data_handle(), + num_levels, + lower_level, + upper_level, + num_vectors, + stream_.get())); // ------------------------- // 4. Compute prefix sum (offsets) on GPU using CUB diff --git a/cpp/src/neighbors/scann/detail/scann_avq.cuh b/cpp/src/neighbors/scann/detail/scann_avq.cuh index d17dcc405f..fb10bedbdf 100644 --- a/cpp/src/neighbors/scann/detail/scann_avq.cuh +++ b/cpp/src/neighbors/scann/detail/scann_avq.cuh @@ -3,6 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "../../../core/cub_histogram.hpp" #include #include #include @@ -68,27 +69,27 @@ void compute_cluster_offsets(raft::resources const& dev_resources, size_t temp_storage_bytes = 0; - cub::DeviceHistogram::HistogramEven(nullptr, - temp_storage_bytes, - clusters.data_handle(), - cluster_sizes.data_handle(), - num_levels, - lower_level, - upper_level, - clusters.extent(0), - stream); + cuvs::detail::histogram_even(nullptr, + temp_storage_bytes, + clusters.data_handle(), + cluster_sizes.data_handle(), + num_levels, + lower_level, + upper_level, + clusters.extent(0), + stream); rmm::device_uvector temp_storage_hist(temp_storage_bytes, stream, device_memory); - cub::DeviceHistogram::HistogramEven(temp_storage_hist.data(), - temp_storage_bytes, - clusters.data_handle(), - cluster_sizes.data_handle(), - num_levels, - lower_level, - upper_level, - clusters.extent(0), - stream); + cuvs::detail::histogram_even(temp_storage_hist.data(), + temp_storage_bytes, + clusters.data_handle(), + cluster_sizes.data_handle(), + num_levels, + lower_level, + upper_level, + clusters.extent(0), + stream); int num_items = cluster_sizes.extent(0); diff --git a/cpp/src/stats/detail/silhouette_score.cuh b/cpp/src/stats/detail/silhouette_score.cuh index 5ba31bd911..4585a843eb 100644 --- a/cpp/src/stats/detail/silhouette_score.cuh +++ b/cpp/src/stats/detail/silhouette_score.cuh @@ -5,6 +5,7 @@ #pragma once +#include "../../core/cub_histogram.hpp" #include #include #include @@ -111,27 +112,27 @@ void countLabels(const LabelT* labels, rmm::device_uvector countArray(nUniqueLabels, stream); - RAFT_CUDA_TRY(cub::DeviceHistogram::HistogramEven(nullptr, - temp_storage_bytes, - labels, - binCountArray, - num_levels, - lower_level, - upper_level, - nRows, - stream)); + RAFT_CUDA_TRY(cuvs::detail::histogram_even(nullptr, + temp_storage_bytes, + labels, + binCountArray, + num_levels, + lower_level, + upper_level, + nRows, + stream)); workspace.resize(temp_storage_bytes, stream); - RAFT_CUDA_TRY(cub::DeviceHistogram::HistogramEven(workspace.data(), - temp_storage_bytes, - labels, - binCountArray, - num_levels, - lower_level, - upper_level, - nRows, - stream)); + RAFT_CUDA_TRY(cuvs::detail::histogram_even(workspace.data(), + temp_storage_bytes, + labels, + binCountArray, + num_levels, + lower_level, + upper_level, + nRows, + stream)); } /** From a6814d2ce111de1a2e9b756a21bd614519c5a0ad Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 27 Oct 2026 21:44:33 +0000 Subject: [PATCH 2/2] Generate one concrete template instantiation per CUDA translation unit --- .pre-commit-config.yaml | 5 + cpp/CMakeLists.txt | 36 +- cpp/src/core/count_labels.cu.in | 4 +- cpp/src/core/count_labels_matrix.json | 12 +- cpp/src/core/cub_histogram_init.cu.in | 19 + cpp/src/core/cub_histogram_init_matrix.json | 62 +++ cpp/src/core/cub_histogram_launchers.hpp | 399 +++++++++++++++++++ cpp/src/core/cub_histogram_sweep.cu.in | 44 ++ cpp/src/core/cub_histogram_sweep_matrix.json | 26 ++ cpp/src/core/histogram_even.cu.in | 5 +- cpp/src/core/histogram_even_matrix.json | 54 ++- 11 files changed, 627 insertions(+), 39 deletions(-) create mode 100644 cpp/src/core/cub_histogram_init.cu.in create mode 100644 cpp/src/core/cub_histogram_init_matrix.json create mode 100644 cpp/src/core/cub_histogram_launchers.hpp create mode 100644 cpp/src/core/cub_histogram_sweep.cu.in create mode 100644 cpp/src/core/cub_histogram_sweep_matrix.json diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 31a707c1dc..d91a944152 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -176,6 +176,7 @@ repos: cpp/cmake/thirdparty/cmake/FindCUDAToolkit[.]cmake$| cpp/cmake/modules/FindAVX[.]cmake$| cpp/src/neighbors/detail/faiss_distance_utils[.]h$| + cpp/src/core/cub_histogram_launchers[.]hpp$| cpp/src/distance/detail/fused_distance_nn/gemm[.]h$| cpp/src/distance/detail/fused_distance_nn/epilogue[.]cuh$| cpp/src/distance/detail/fused_distance_nn/predicated_tile_iterator_normvec_smem[.]h$| @@ -185,6 +186,10 @@ repos: cpp/src/distance/detail/fused_distance_nn/custom_epilogue_with_broadcast[.]h$| rust/cuvs-sys/src/bindings[.]rs$ ) + - id: verify-copyright + name: verify-copyright-cub + args: [--fix, --spdx, "--spdx-license-identifier=Apache-2.0 AND BSD-3-Clause"] + files: ^cpp/src/core/cub_histogram_launchers[.]hpp$ - id: verify-copyright name: verify-copyright-cmake args: [--fix, --spdx, "--spdx-license-identifier=Apache-2.0 AND BSD-3-Clause"] diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 6baaf073d6..4e4521ca6f 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -1431,25 +1431,31 @@ if(NOT BUILD_CPU_ONLY) "${CMAKE_CURRENT_BINARY_DIR}/src/selection/select_k_inst_data_@data_abbrev@_index_@index_abbrev@.cu" ) - # Keep every primitive function in its own CUDA translation unit. foreach(primitive IN ITEMS src/core/count_labels src/core/histogram_even) - set(primitive_instantiations) - generate_string_matrix( - primitive_instantiations - MATRIX_JSON_FILE - "${CMAKE_CURRENT_SOURCE_DIR}/${primitive}_matrix.json" - ITEM_FORMAT - "template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@@semicolon@" - GLUE - "\n" - ) - configure_file( - "${CMAKE_CURRENT_SOURCE_DIR}/${primitive}.cu.in" - "${CMAKE_CURRENT_BINARY_DIR}/${primitive}.cu" @ONLY + generate_inst_matrix( + cuvs_primitive_sources + MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/${primitive}_matrix.json" + INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/${primitive}.cu.in" + OUTPUT_FILE_FORMAT "${CMAKE_CURRENT_BINARY_DIR}/${primitive}_@instantiation_name@.cu" ) - list(APPEND cuvs_primitive_sources "${CMAKE_CURRENT_BINARY_DIR}/${primitive}.cu") endforeach() + generate_inst_matrix( + cuvs_primitive_sources + MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/cub_histogram_init_matrix.json" + INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/cub_histogram_init.cu.in" + OUTPUT_FILE_FORMAT + "${CMAKE_CURRENT_BINARY_DIR}/src/core/cub_histogram_init_@instantiation_name@.cu" + ) + + generate_inst_matrix( + cuvs_primitive_sources + MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/cub_histogram_sweep_matrix.json" + INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/cub_histogram_sweep.cu.in" + OUTPUT_FILE_FORMAT + "${CMAKE_CURRENT_BINARY_DIR}/src/core/cub_histogram_sweep_@instantiation_name@.cu" + ) + add_library( cuvs_objs OBJECT ${cuvs_primitive_sources} diff --git a/cpp/src/core/count_labels.cu.in b/cpp/src/core/count_labels.cu.in index 532bc84761..22fd3bc1bc 100644 --- a/cpp/src/core/count_labels.cu.in +++ b/cpp/src/core/count_labels.cu.in @@ -44,6 +44,8 @@ void count_labels(const uint32_t* labels, stream)); } -@primitive_instantiations@ +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on } // namespace cuvs::detail diff --git a/cpp/src/core/count_labels_matrix.json b/cpp/src/core/count_labels_matrix.json index 78fdd81656..d0db44a483 100644 --- a/cpp/src/core/count_labels_matrix.json +++ b/cpp/src/core/count_labels_matrix.json @@ -7,19 +7,23 @@ "_variants": [ { "template_arguments": "", - "parameter_types": "const uint32_t*, uint32_t*, int, int, rmm::device_uvector&, cudaStream_t" + "parameter_types": "const uint32_t*, uint32_t*, int, int, rmm::device_uvector&, cudaStream_t", + "instantiation_name": "u32_int" }, { "template_arguments": "", - "parameter_types": "const uint32_t*, uint32_t*, int64_t, int64_t, rmm::device_uvector&, cudaStream_t" + "parameter_types": "const uint32_t*, uint32_t*, int64_t, int64_t, rmm::device_uvector&, cudaStream_t", + "instantiation_name": "u32_i64" }, { "template_arguments": "", - "parameter_types": "const uint32_t*, unsigned long long*, int, int, rmm::device_uvector&, cudaStream_t" + "parameter_types": "const uint32_t*, unsigned long long*, int, int, rmm::device_uvector&, cudaStream_t", + "instantiation_name": "u64_int" }, { "template_arguments": "", - "parameter_types": "const uint32_t*, unsigned long long*, int64_t, int64_t, rmm::device_uvector&, cudaStream_t" + "parameter_types": "const uint32_t*, unsigned long long*, int64_t, int64_t, rmm::device_uvector&, cudaStream_t", + "instantiation_name": "u64_i64" } ] } diff --git a/cpp/src/core/cub_histogram_init.cu.in b/cpp/src/core/cub_histogram_init.cu.in new file mode 100644 index 0000000000..22bac8130e --- /dev/null +++ b/cpp/src/core/cub_histogram_init.cu.in @@ -0,0 +1,19 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include + +namespace cuvs::detail::histogram_launch { +template +CUB_RUNTIME_FUNCTION init_pointer_t init() +{ + return &cub::detail::histogram::DeviceHistogramInitKernel; +} +} // namespace cuvs::detail::histogram_launch + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(); +// clang-format on diff --git a/cpp/src/core/cub_histogram_init_matrix.json b/cpp/src/core/cub_histogram_init_matrix.json new file mode 100644 index 0000000000..bd77e347a5 --- /dev/null +++ b/cpp/src/core/cub_histogram_init_matrix.json @@ -0,0 +1,62 @@ +[ + { + "return_type": "cuvs::detail::histogram_launch::init_pointer_t, 1, unsigned long long, int>", + "qualified_function": "cuvs::detail::histogram_launch::init", + "template_arguments": ", 1, unsigned long long, int>", + "instantiation_name": "u32_count_u64_offset_i32" + }, + { + "return_type": "cuvs::detail::histogram_launch::init_pointer_t, 1, float, int>", + "qualified_function": "cuvs::detail::histogram_launch::init", + "template_arguments": ", 1, float, int>", + "instantiation_name": "i32_count_f32_offset_i32" + }, + { + "return_type": "cuvs::detail::histogram_launch::init_pointer_t, 1, double, long>", + "qualified_function": "cuvs::detail::histogram_launch::init", + "template_arguments": ", 1, double, long>", + "instantiation_name": "i64_count_f64_offset_i64" + }, + { + "return_type": "cuvs::detail::histogram_launch::init_pointer_t, 1, double, int>", + "qualified_function": "cuvs::detail::histogram_launch::init", + "template_arguments": ", 1, double, int>", + "instantiation_name": "i64_count_f64_offset_i32" + }, + { + "return_type": "cuvs::detail::histogram_launch::init_pointer_t, 1, float, long>", + "qualified_function": "cuvs::detail::histogram_launch::init", + "template_arguments": ", 1, float, long>", + "instantiation_name": "i64_count_f32_offset_i64" + }, + { + "return_type": "cuvs::detail::histogram_launch::init_pointer_t, 1, float, int>", + "qualified_function": "cuvs::detail::histogram_launch::init", + "template_arguments": ", 1, float, int>", + "instantiation_name": "i64_count_f32_offset_i32" + }, + { + "return_type": "cuvs::detail::histogram_launch::init_pointer_t, 1, unsigned int, int>", + "qualified_function": "cuvs::detail::histogram_launch::init", + "template_arguments": ", 1, unsigned int, int>", + "instantiation_name": "u32_count_u32_offset_i32" + }, + { + "return_type": "cuvs::detail::histogram_launch::init_pointer_t, 1, double, int>", + "qualified_function": "cuvs::detail::histogram_launch::init", + "template_arguments": ", 1, double, int>", + "instantiation_name": "i32_count_f64_offset_i32" + }, + { + "return_type": "cuvs::detail::histogram_launch::init_pointer_t, 1, int, int>", + "qualified_function": "cuvs::detail::histogram_launch::init", + "template_arguments": ", 1, int, int>", + "instantiation_name": "i32_count_i32_offset_i32" + }, + { + "return_type": "cuvs::detail::histogram_launch::init_pointer_t, 1, unsigned int, long>", + "qualified_function": "cuvs::detail::histogram_launch::init", + "template_arguments": ", 1, unsigned int, long>", + "instantiation_name": "u32_count_u32_offset_i64" + } +] diff --git a/cpp/src/core/cub_histogram_launchers.hpp b/cpp/src/core/cub_histogram_launchers.hpp new file mode 100644 index 0000000000..fb3bfedd37 --- /dev/null +++ b/cpp/src/core/cub_histogram_launchers.hpp @@ -0,0 +1,399 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2011, Duane Merrill. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2011-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 AND BSD-3-Clause + */ + +#pragma once + +#include +#include +#include + +namespace cuvs::detail::histogram_launch { + +template +using init_pointer_t = void (*)(cuda::std::array, + cuda::std::array, + cub::GridQueue); + +template +CUB_RUNTIME_FUNCTION init_pointer_t init(); + +template +using sweep_pointer_t = void (*)(Samples, + cuda::std::array, + cuda::std::array, + cuda::std::array, + cuda::std::array, + cuda::std::array, + cuda::std::array, + Offset, + Offset, + Offset, + int, + cub::GridQueue); + +template +CUB_RUNTIME_FUNCTION sweep_pointer_t +sweep(); + +// Adapted from CUB DeviceHistogramKernelSource; kernel implementations remain in CUB. +template +struct source { + using TransformsT = cub::detail::histogram::Transforms; + + template + _CCCL_HIDE_FROM_ABI CUB_RUNTIME_FUNCTION static constexpr auto HistogramInitKernel() + { + return init(); + } + + /// Returns the default histogram sweep kernel that receives pre-initialized decode operators from + /// the host. + template + _CCCL_HIDE_FROM_ABI CUB_RUNTIME_FUNCTION static constexpr auto HistogramSweepKernel() + { + if constexpr (::cuda::std::is_same_v && + ::cuda::std::is_same_v && ::cuda::std::is_same_v && + (::cuda::std::is_same_v || + ::cuda::std::is_same_v)) { + return sweep(); + } else { + return &cub::detail::histogram::DeviceHistogramSweepKernel; + } + } + + /// Returns the device-init histogram sweep kernel that initializes decode operators from level + /// arrays in the kernel. + template + _CCCL_HIDE_FROM_ABI CUB_RUNTIME_FUNCTION static constexpr auto HistogramSweepKernelDeviceInit() + { + // For DispatchEven, we use the scale transform to convert samples to + // privatized bins and pass-thru transform to convert privatized bins to + // output bins, vice verse for byte samples. + + // For DispatchRange, we use the search transform to convert samples to + // privatized bins and scale transform to convert privatized bins to output bins, + // vice verse for byte samples. + + using DecodeOpT = + ::cuda::std::conditional_t>; + + using PrivatizedDecodeOpT = + ::cuda::std::conditional_t; + using OutputDecodeOpT = + ::cuda::std::conditional_t; + + return &cub::detail::histogram::DeviceHistogramSweepDeviceInitKernel; + } + + CUB_RUNTIME_FUNCTION static constexpr size_t CounterSize() { return sizeof(CounterT); } + + template + CUB_RUNTIME_FUNCTION static constexpr bool MayOverflow( + [[maybe_unused]] NumBinsT num_bins, + [[maybe_unused]] const UpperLevelArrayT& upper_level, + [[maybe_unused]] const LowerLevelArrayT& lower_level, + [[maybe_unused]] int channel) + { + using CommonT = typename TransformsT::ScaleTransform::CommonT; + + if constexpr (::cuda::std::is_integral_v) { + using IntArithmeticT = typename TransformsT::ScaleTransform::IntArithmeticT; + // The unary plus promotes plain char to int, which cuda::std::cmp_greater requires + if (::cuda::std::cmp_greater(num_bins, +::cuda::std::numeric_limits::max())) { + return true; + } + const auto upper_level_cast = static_cast(upper_level[channel]); + const auto lower_level_cast = static_cast(lower_level[channel]); + const auto range = static_cast(upper_level_cast - lower_level_cast); + return range > (::cuda::std::numeric_limits::max() / + static_cast(num_bins)); + } else { + return false; + } + } +}; + +} // namespace cuvs::detail::histogram_launch + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + double, + int, + int, + int> + : cuvs::detail::histogram_launch::source< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + double, + int, + int, + int> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + float, + int, + int, + int> + : cuvs::detail::histogram_launch::source< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + float, + int, + int, + int> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + double, + long, + int, + long> + : cuvs::detail::histogram_launch::source< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + double, + long, + int, + long> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + double, + long, + long, + long> + : cuvs::detail::histogram_launch::source< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + double, + long, + long, + long> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + float, + long, + int, + long> + : cuvs::detail::histogram_launch::source< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + float, + long, + int, + long> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + float, + long, + long, + long> + : cuvs::detail::histogram_launch::source< + 1, + 1, + ::cuda::transform_iterator< + cuvs::distance::detail::Top1nnResultTraits*>::key_op, + raft::KeyValuePair*>, + float, + long, + long, + long> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, 1, int const*, double, int, int, int> + : cuvs::detail::histogram_launch::source<1, 1, int const*, double, int, int, int> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, 1, int const*, float, int, int, int> + : cuvs::detail::histogram_launch::source<1, 1, int const*, float, int, int, int> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, 1, int const*, int, int, int, int> + : cuvs::detail::histogram_launch::source<1, 1, int const*, int, int, int, int> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, 1, int*, double, int, int, int> + : cuvs::detail::histogram_launch::source<1, 1, int*, double, int, int, int> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, 1, int*, float, int, int, int> + : cuvs::detail::histogram_launch::source<1, 1, int*, float, int, int, int> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, 1, int*, int, int, int, int> + : cuvs::detail::histogram_launch::source<1, 1, int*, int, int, int, int> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, 1, long*, double, long, int, long> + : cuvs::detail::histogram_launch::source<1, 1, long*, double, long, int, long> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, 1, long*, double, long, long, long> + : cuvs::detail::histogram_launch::source<1, 1, long*, double, long, long, long> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, 1, long*, float, long, int, long> + : cuvs::detail::histogram_launch::source<1, 1, long*, float, long, int, long> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, 1, long*, float, long, long, long> + : cuvs::detail::histogram_launch::source<1, 1, long*, float, long, long, long> {}; + +template <> +struct cub::detail::histogram:: + DeviceHistogramKernelSource<1, 1, unsigned int const*, unsigned int, int, int, unsigned int> + : cuvs::detail::histogram_launch:: + source<1, 1, unsigned int const*, unsigned int, int, int, unsigned int> {}; + +template <> +struct cub::detail::histogram:: + DeviceHistogramKernelSource<1, 1, unsigned int const*, unsigned int, int, long, unsigned int> + : cuvs::detail::histogram_launch:: + source<1, 1, unsigned int const*, unsigned int, int, long, unsigned int> {}; + +template <> +struct cub::detail::histogram:: + DeviceHistogramKernelSource<1, 1, unsigned int const*, unsigned long long, int, int, unsigned int> + : cuvs::detail::histogram_launch:: + source<1, 1, unsigned int const*, unsigned long long, int, int, unsigned int> {}; + +template <> +struct cub::detail::histogram:: + DeviceHistogramKernelSource<1, 1, unsigned int const*, unsigned int, long, int, unsigned int> + : cuvs::detail::histogram_launch:: + source<1, 1, unsigned int const*, unsigned int, long, int, unsigned int> {}; + +template <> +struct cub::detail::histogram:: + DeviceHistogramKernelSource<1, 1, unsigned int const*, unsigned int, long, long, unsigned int> + : cuvs::detail::histogram_launch:: + source<1, 1, unsigned int const*, unsigned int, long, long, unsigned int> {}; + +template <> +struct cub::detail::histogram::DeviceHistogramKernelSource<1, + 1, + unsigned int const*, + unsigned long long, + long, + int, + unsigned int> + : cuvs::detail::histogram_launch:: + source<1, 1, unsigned int const*, unsigned long long, long, int, unsigned int> {}; diff --git a/cpp/src/core/cub_histogram_sweep.cu.in b/cpp/src/core/cub_histogram_sweep.cu.in new file mode 100644 index 0000000000..4fbdd2f5b4 --- /dev/null +++ b/cpp/src/core/cub_histogram_sweep.cu.in @@ -0,0 +1,44 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include + +namespace cuvs::detail::histogram_launch { +template +CUB_RUNTIME_FUNCTION sweep_pointer_t +sweep() +{ + return &cub::detail::histogram::DeviceHistogramSweepKernel; +} +} // namespace cuvs::detail::histogram_launch + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(); +// clang-format on diff --git a/cpp/src/core/cub_histogram_sweep_matrix.json b/cpp/src/core/cub_histogram_sweep_matrix.json new file mode 100644 index 0000000000..b5e4235518 --- /dev/null +++ b/cpp/src/core/cub_histogram_sweep_matrix.json @@ -0,0 +1,26 @@ +[ + { + "return_type": "cuvs::detail::histogram_launch::sweep_pointer_t, 256, 1, 1, unsigned int const*, unsigned long long, cub::detail::histogram::Transforms::ScaleTransform, cub::detail::histogram::Transforms::PassThruTransform, int>", + "qualified_function": "cuvs::detail::histogram_launch::sweep", + "template_arguments": ", 256, 1, 1, unsigned int const*, unsigned long long, cub::detail::histogram::Transforms::ScaleTransform, cub::detail::histogram::Transforms::PassThruTransform, int>", + "instantiation_name": "u32_count_u64_offset_i32_bins_256" + }, + { + "return_type": "cuvs::detail::histogram_launch::sweep_pointer_t, 0, 1, 1, unsigned int const*, unsigned long long, cub::detail::histogram::Transforms::ScaleTransform, cub::detail::histogram::Transforms::PassThruTransform, int>", + "qualified_function": "cuvs::detail::histogram_launch::sweep", + "template_arguments": ", 0, 1, 1, unsigned int const*, unsigned long long, cub::detail::histogram::Transforms::ScaleTransform, cub::detail::histogram::Transforms::PassThruTransform, int>", + "instantiation_name": "u32_count_u64_offset_i32_bins_0" + }, + { + "return_type": "cuvs::detail::histogram_launch::sweep_pointer_t, 256, 1, 1, unsigned int const*, unsigned int, cub::detail::histogram::Transforms::ScaleTransform, cub::detail::histogram::Transforms::PassThruTransform, int>", + "qualified_function": "cuvs::detail::histogram_launch::sweep", + "template_arguments": ", 256, 1, 1, unsigned int const*, unsigned int, cub::detail::histogram::Transforms::ScaleTransform, cub::detail::histogram::Transforms::PassThruTransform, int>", + "instantiation_name": "u32_count_u32_offset_i32_bins_256" + }, + { + "return_type": "cuvs::detail::histogram_launch::sweep_pointer_t, 0, 1, 1, unsigned int const*, unsigned int, cub::detail::histogram::Transforms::ScaleTransform, cub::detail::histogram::Transforms::PassThruTransform, int>", + "qualified_function": "cuvs::detail::histogram_launch::sweep", + "template_arguments": ", 0, 1, 1, unsigned int const*, unsigned int, cub::detail::histogram::Transforms::ScaleTransform, cub::detail::histogram::Transforms::PassThruTransform, int>", + "instantiation_name": "u32_count_u32_offset_i32_bins_0" + } +] diff --git a/cpp/src/core/histogram_even.cu.in b/cpp/src/core/histogram_even.cu.in index 698a501816..da6e084fda 100644 --- a/cpp/src/core/histogram_even.cu.in +++ b/cpp/src/core/histogram_even.cu.in @@ -3,6 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include #include #include @@ -29,6 +30,8 @@ cudaError_t histogram_even(void* storage, storage, bytes, samples, counts, levels, lower, upper, count, stream); } -@primitive_instantiations@ +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on } // namespace cuvs::detail diff --git a/cpp/src/core/histogram_even_matrix.json b/cpp/src/core/histogram_even_matrix.json index 5dcb79543f..933384d0e2 100644 --- a/cpp/src/core/histogram_even_matrix.json +++ b/cpp/src/core/histogram_even_matrix.json @@ -7,75 +7,93 @@ "_variants": [ { "template_arguments": "< ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, double, int, int>", - "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, double*, int, int, int, int, cudaStream_t" + "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, double*, int, int, int, int, cudaStream_t", + "instantiation_name": "transform_top1nn_key_value_int_f64_ptr_key_op_key_value_int_f64_ptr_f64_int_int" }, { "template_arguments": "< ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, float, int, int>", - "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, float*, int, int, int, int, cudaStream_t" + "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, float*, int, int, int, int, cudaStream_t", + "instantiation_name": "transform_top1nn_key_value_int_f32_ptr_key_op_key_value_int_f32_ptr_f32_int_int" }, { "template_arguments": "< ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, double, long, long>", - "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, double*, int, long, long, long, cudaStream_t" + "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, double*, int, long, long, long, cudaStream_t", + "instantiation_name": "transform_top1nn_key_value_long_f64_ptr_key_op_key_value_long_f64_ptr_f64_long_long" }, { "template_arguments": "< ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, float, long, long>", - "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, float*, int, long, long, long, cudaStream_t" + "parameter_types": "void*, unsigned long&, ::cuda::transform_iterator< cuvs::distance::detail::Top1nnResultTraits*>::key_op, raft::KeyValuePair*>, float*, int, long, long, long, cudaStream_t", + "instantiation_name": "transform_top1nn_key_value_long_f32_ptr_key_op_key_value_long_f32_ptr_f32_long_long" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, int const*, double*, int, int, int, int, cudaStream_t" + "parameter_types": "void*, unsigned long&, int const*, double*, int, int, int, int, cudaStream_t", + "instantiation_name": "int_c_ptr_f64_int_int" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, int const*, float*, int, int, int, int, cudaStream_t" + "parameter_types": "void*, unsigned long&, int const*, float*, int, int, int, int, cudaStream_t", + "instantiation_name": "int_c_ptr_f32_int_int" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, int const*, int*, int, int, int, int, cudaStream_t" + "parameter_types": "void*, unsigned long&, int const*, int*, int, int, int, int, cudaStream_t", + "instantiation_name": "int_c_ptr_int_int_int" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, int*, double*, int, int, int, int, cudaStream_t" + "parameter_types": "void*, unsigned long&, int*, double*, int, int, int, int, cudaStream_t", + "instantiation_name": "int_ptr_f64_int_int" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, int*, float*, int, int, int, int, cudaStream_t" + "parameter_types": "void*, unsigned long&, int*, float*, int, int, int, int, cudaStream_t", + "instantiation_name": "int_ptr_f32_int_int" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, int*, int*, int, int, int, int, cudaStream_t" + "parameter_types": "void*, unsigned long&, int*, int*, int, int, int, int, cudaStream_t", + "instantiation_name": "int_ptr_int_int_int" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, long*, double*, int, long, long, long, cudaStream_t" + "parameter_types": "void*, unsigned long&, long*, double*, int, long, long, long, cudaStream_t", + "instantiation_name": "long_ptr_f64_long_long" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, long*, float*, int, long, long, long, cudaStream_t" + "parameter_types": "void*, unsigned long&, long*, float*, int, long, long, long, cudaStream_t", + "instantiation_name": "long_ptr_f32_long_long" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned int*, int, int, int, long, cudaStream_t" + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned int*, int, int, int, long, cudaStream_t", + "instantiation_name": "u32_c_ptr_u32_int_long" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned long long*, int, int, int, unsigned long, cudaStream_t" + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned long long*, int, int, int, unsigned long, cudaStream_t", + "instantiation_name": "u32_c_ptr_u64_int_ulong" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned int*, int, int, int, int, cudaStream_t" + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned int*, int, int, int, int, cudaStream_t", + "instantiation_name": "u32_c_ptr_u32_int_int" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned int*, int, long, long, long, cudaStream_t" + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned int*, int, long, long, long, cudaStream_t", + "instantiation_name": "u32_c_ptr_u32_long_long" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned long long*, int, int, int, int, cudaStream_t" + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned long long*, int, int, int, int, cudaStream_t", + "instantiation_name": "u32_c_ptr_u64_int_int" }, { "template_arguments": "", - "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned long long*, int, long, long, long, cudaStream_t" + "parameter_types": "void*, unsigned long&, unsigned int const*, unsigned long long*, int, long, long, long, cudaStream_t", + "instantiation_name": "u32_c_ptr_u64_long_long" } ] }