diff --git a/cpp/cmake/thirdparty/get_raft.cmake b/cpp/cmake/thirdparty/get_raft.cmake index 6a08edb2ab..60bf692917 100644 --- a/cpp/cmake/thirdparty/get_raft.cmake +++ b/cpp/cmake/thirdparty/get_raft.cmake @@ -13,8 +13,8 @@ # Use RAPIDS_VERSION_MAJOR_MINOR from rapids_config.cmake set(RAFT_VERSION "${RAPIDS_VERSION_MAJOR_MINOR}") -set(RAFT_FORK "rapidsai") -set(RAFT_PINNED_TAG "${rapids-cmake-checkout-tag}") +set(RAFT_FORK "bdice") +set(RAFT_PINNED_TAG "cccl-mdspan") function(find_and_configure_raft) set(oneValueArgs VERSION FORK PINNED_TAG BUILD_STATIC_DEPS ENABLE_NVTX ENABLE_MNMG_DEPENDENCIES CLONE_ON_PIN) diff --git a/cpp/include/cuvs/cluster/kmeans.hpp b/cpp/include/cuvs/cluster/kmeans.hpp index 43d2db5a5c..0de350c5a9 100644 --- a/cpp/include/cuvs/cluster/kmeans.hpp +++ b/cpp/include/cuvs/cluster/kmeans.hpp @@ -1233,7 +1233,7 @@ namespace helpers { * * #include * - * using namespace cuvs::cluster; + * using namespace cuvs::cluster; * * raft::handle_t handle; * int n_samples = 100, n_features = 15, n_clusters = 10; @@ -1242,9 +1242,9 @@ namespace helpers { * * raft::random::make_blobs(handle, X, labels, n_clusters); * - * auto best_k = raft::make_host_scalar(0); - * auto n_iter = raft::make_host_scalar(0); - * auto inertia = raft::make_host_scalar(0); + * auto best_k = raft::make_host_scalar(0); + * auto n_iter = raft::make_host_scalar(0); + * auto inertia = raft::make_host_scalar(0); * * kmeans::find_k(handle, X, best_k.view(), inertia.view(), n_iter.view(), n_clusters+1); * diff --git a/cpp/include/cuvs/neighbors/common.hpp b/cpp/include/cuvs/neighbors/common.hpp index 4ed79dd695..089925382a 100644 --- a/cpp/include/cuvs/neighbors/common.hpp +++ b/cpp/include/cuvs/neighbors/common.hpp @@ -259,7 +259,7 @@ auto make_strided_dataset(const raft::resources& res, const SrcT& src, uint32_t } // Something is wrong: have to make a copy and produce an owning dataset auto out_layout = - raft::make_strided_layout(src.extents(), std::array{required_stride, 1}); + raft::make_strided_layout(src.extents(), cuda::std::array{required_stride, 1}); auto out_array = raft::make_device_matrix(res, src.extent(0), required_stride); @@ -321,7 +321,7 @@ auto make_strided_dataset( const bool stride_matches = required_stride == src_stride; auto out_layout = - raft::make_strided_layout(src.extents(), std::array{required_stride, 1}); + raft::make_strided_layout(src.extents(), cuda::std::array{required_stride, 1}); using out_mdarray_type = raft::device_matrix; using out_layout_type = typename out_mdarray_type::layout_type; diff --git a/cpp/include/cuvs/neighbors/ivf_pq.hpp b/cpp/include/cuvs/neighbors/ivf_pq.hpp index 593403f2aa..9dd8e8fb6d 100644 --- a/cpp/include/cuvs/neighbors/ivf_pq.hpp +++ b/cpp/include/cuvs/neighbors/ivf_pq.hpp @@ -346,8 +346,8 @@ struct index : cuvs::neighbors::index { static_assert(!raft::is_narrowing_v, "IdxT must be able to represent all values of uint32_t"); - using pq_centers_extents = std::experimental:: - extents; + using pq_centers_extents = + cuda::std::extents; public: index(const index&) = delete; diff --git a/cpp/src/cluster/detail/kmeans.cuh b/cpp/src/cluster/detail/kmeans.cuh index d5808d6480..9b50ff55c8 100644 --- a/cpp/src/cluster/detail/kmeans.cuh +++ b/cpp/src/cluster/detail/kmeans.cuh @@ -355,8 +355,8 @@ void kmeans_fit_main(raft::resources const& handle, raft::device_matrix_view X, raft::device_vector_view weight, raft::device_matrix_view centroidsRawData, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter, + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter, rmm::device_uvector& workspace) { raft::common::nvtx::range fun_scope("kmeans_fit_main"); @@ -738,8 +738,8 @@ void initScalableKMeansPlusPlus(raft::resources const& handle, cuvs::cluster::kmeans::detail::kmeansPlusPlus( handle, params, potentialCentroids, centroidsRawData, workspace); - auto inertia = raft::make_host_scalar(0); - auto n_iter = raft::make_host_scalar(0); + auto inertia = raft::make_host_scalar(0); + auto n_iter = raft::make_host_scalar(0); cuvs::cluster::kmeans::params default_params; default_params.n_clusters = params.n_clusters; @@ -814,8 +814,8 @@ void kmeans_fit(raft::resources const& handle, raft::device_matrix_view X, std::optional> sample_weight, raft::device_matrix_view centroids, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter) + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter) { raft::common::nvtx::range fun_scope("kmeans_fit"); auto n_samples = X.extent(0); @@ -859,7 +859,7 @@ void kmeans_fit(raft::resources const& handle, // Allocate memory rmm::device_uvector workspace(0, stream); - auto weight = raft::make_device_vector(handle, n_samples); + auto weight = raft::make_device_vector(handle, n_samples); if (sample_weight.has_value()) raft::copy(weight.data_handle(), sample_weight.value().data_handle(), n_samples, stream); else @@ -907,12 +907,13 @@ void kmeans_fit(raft::resources const& handle, "k-means++ algorithm.", seed_iter + 1, n_init); - if (iter_params.oversampling_factor == 0) + if (iter_params.oversampling_factor == 0) { cuvs::cluster::kmeans::detail::kmeansPlusPlus( handle, iter_params, X, centroidsRawData.view(), workspace); - else + } else { cuvs::cluster::kmeans::detail::initScalableKMeansPlusPlus( handle, iter_params, X, centroidsRawData.view(), workspace); + } } else if (iter_params.init == cuvs::cluster::kmeans::params::InitMethod::Array) { RAFT_LOG_DEBUG( "KMeans.fit (Iteration-%d/%d): initialize cluster centers from " @@ -932,8 +933,8 @@ void kmeans_fit(raft::resources const& handle, X, weight.view(), centroidsRawData.view(), - raft::make_host_scalar_view(&iter_inertia), - raft::make_host_scalar_view(&n_current_iter), + raft::make_host_scalar_view(&iter_inertia), + raft::make_host_scalar_view(&n_current_iter), workspace); if (iter_inertia < inertia[0]) { inertia[0] = iter_inertia; @@ -964,12 +965,12 @@ void kmeans_fit(raft::resources const& handle, auto XView = raft::make_device_matrix_view(X, n_samples, n_features); auto centroidsView = raft::make_device_matrix_view(centroids, pams.n_clusters, n_features); - std::optional> sample_weightView = std::nullopt; + std::optional> sample_weightView = std::nullopt; if (sample_weight) sample_weightView = raft::make_device_vector_view(sample_weight, n_samples); - auto inertiaView = raft::make_host_scalar_view(&inertia); - auto n_iterView = raft::make_host_scalar_view(&n_iter); + auto inertiaView = raft::make_host_scalar_view(&inertia); + auto n_iterView = raft::make_host_scalar_view(&n_iter); cuvs::cluster::kmeans::detail::kmeans_fit( handle, pams, XView, sample_weightView, centroidsView, inertiaView, n_iterView); @@ -983,7 +984,7 @@ void kmeans_predict(raft::resources const& handle, raft::device_matrix_view centroids, raft::device_vector_view labels, bool normalize_weight, - raft::host_scalar_view inertia) + raft::host_scalar_view inertia) { raft::common::nvtx::range fun_scope("kmeans_predict"); auto n_samples = X.extent(0); @@ -1102,7 +1103,7 @@ void kmeans_predict(raft::resources const& handle, sample_weightView.emplace( raft::make_device_vector_view(sample_weight, n_samples)); auto labelsView = raft::make_device_vector_view(labels, n_samples); - auto inertiaView = raft::make_host_scalar_view(&inertia); + auto inertiaView = raft::make_host_scalar_view(&inertia); cuvs::cluster::kmeans::detail::kmeans_predict(handle, pams, @@ -1121,8 +1122,8 @@ void kmeans_fit_predict(raft::resources const& handle, std::optional> sample_weight, std::optional> centroids, raft::device_vector_view labels, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter) + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter) { raft::common::nvtx::range fun_scope("kmeans_fit_predict"); if (!centroids.has_value()) { @@ -1163,8 +1164,8 @@ void kmeans_fit_predict(raft::resources const& handle, centroidsView.emplace( raft::make_device_matrix_view(centroids, pams.n_clusters, n_features)); auto labelsView = raft::make_device_vector_view(labels, n_samples); - auto inertiaView = raft::make_host_scalar_view(&inertia); - auto n_iterView = raft::make_host_scalar_view(&n_iter); + auto inertiaView = raft::make_host_scalar_view(&inertia); + auto n_iterView = raft::make_host_scalar_view(&n_iter); cuvs::cluster::kmeans::detail::kmeans_fit_predict( handle, pams, XView, sample_weightView, centroidsView, labelsView, inertiaView, n_iterView); @@ -1184,9 +1185,9 @@ void kmeans_fit_predict(raft::resources const& handle, template void kmeans_transform(raft::resources const& handle, const cuvs::cluster::kmeans::params& pams, - raft::device_matrix_view X, - raft::device_matrix_view centroids, - raft::device_matrix_view X_new) + raft::device_matrix_view X, + raft::device_matrix_view centroids, + raft::device_matrix_view X_new) { raft::common::nvtx::range fun_scope("kmeans_transform"); raft::default_logger().set_level(pams.verbosity); diff --git a/cpp/src/cluster/detail/kmeans_auto_find_k.cuh b/cpp/src/cluster/detail/kmeans_auto_find_k.cuh index 797b33bca9..e258fe19b2 100644 --- a/cpp/src/cluster/detail/kmeans_auto_find_k.cuh +++ b/cpp/src/cluster/detail/kmeans_auto_find_k.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -35,13 +35,13 @@ void compute_dispersion(raft::resources const& handle, raft::device_matrix_view X, cuvs::cluster::kmeans::params& params, raft::device_matrix_view centroids_view, - raft::device_vector_view labels, - raft::device_vector_view clusterSizes, + raft::device_vector_view labels, + raft::device_vector_view clusterSizes, rmm::device_uvector& workspace, - raft::host_vector_view clusterDispertionView, - raft::host_vector_view resultsView, - raft::host_scalar_view residual, - raft::host_scalar_view n_iter, + raft::host_vector_view clusterDispertionView, + raft::host_vector_view resultsView, + raft::host_scalar_view residual, + raft::host_scalar_view n_iter, int val, idx_t n, idx_t d) @@ -68,9 +68,9 @@ void compute_dispersion(raft::resources const& handle, template void find_k(raft::resources const& handle, raft::device_matrix_view X, - raft::host_scalar_view best_k, - raft::host_scalar_view residual, - raft::host_scalar_view n_iter, + raft::host_scalar_view best_k, + raft::host_scalar_view residual, + raft::host_scalar_view n_iter, idx_t kmax, idx_t kmin = 1, idx_t maxiter = 100, @@ -89,16 +89,16 @@ void find_k(raft::resources const& handle, // Device memory auto centroids = raft::make_device_matrix(handle, kmax, X.extent(1)); - auto clusterSizes = raft::make_device_vector(handle, kmax); - auto labels = raft::make_device_vector(handle, n); + auto clusterSizes = raft::make_device_vector(handle, kmax); + auto labels = raft::make_device_vector(handle, n); rmm::device_uvector workspace(0, raft::resource::get_cuda_stream(handle)); idx_t* clusterSizes_ptr = clusterSizes.data_handle(); // Host memory - auto results = raft::make_host_vector(kmax + 1); - auto clusterDispersion = raft::make_host_vector(kmax + 1); + auto results = raft::make_host_vector(kmax + 1); + auto clusterDispersion = raft::make_host_vector(kmax + 1); auto clusterDispertionView = clusterDispersion.view(); auto resultsView = results.view(); diff --git a/cpp/src/cluster/detail/kmeans_mg.cuh b/cpp/src/cluster/detail/kmeans_mg.cuh index 797e531f58..e454fda2d5 100644 --- a/cpp/src/cluster/detail/kmeans_mg.cuh +++ b/cpp/src/cluster/detail/kmeans_mg.cuh @@ -420,8 +420,8 @@ void initKMeansPlusPlus(const raft::resources& handle, cuvs::cluster::kmeans::init_plus_plus( handle, params, const_centroids, centroidsRawData, workspace); - auto inertia = raft::make_host_scalar(0); - auto n_iter = raft::make_host_scalar(0); + auto inertia = raft::make_host_scalar(0); + auto n_iter = raft::make_host_scalar(0); auto weight_view = raft::make_device_vector_view(weight.data_handle(), weight.extent(0)); cuvs::cluster::kmeans::params params_copy = params; diff --git a/cpp/src/cluster/detail/single_linkage.cuh b/cpp/src/cluster/detail/single_linkage.cuh index e6d00b5b0f..f1d39bfab7 100644 --- a/cpp/src/cluster/detail/single_linkage.cuh +++ b/cpp/src/cluster/detail/single_linkage.cuh @@ -54,7 +54,7 @@ namespace cuvs::cluster::agglomerative::detail { template >> + typename Accessor = raft::device_accessor>> void build_mr_linkage( raft::resources const& handle, raft::mdspan, raft::row_major, Accessor> X, diff --git a/cpp/src/cluster/kmeans.cuh b/cpp/src/cluster/kmeans.cuh index ae9703081a..943eba5531 100644 --- a/cpp/src/cluster/kmeans.cuh +++ b/cpp/src/cluster/kmeans.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2024, NVIDIA CORPORATION. + * Copyright (c) 2020-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -93,8 +93,8 @@ void fit(raft::resources const& handle, raft::device_matrix_view X, std::optional> sample_weight, raft::device_matrix_view centroids, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter) + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter) { // use the mnmg kmeans fit if we have comms initialize, single gpu otherwise if (raft::resource::comms_initialized(handle)) { @@ -165,7 +165,7 @@ void predict(raft::resources const& handle, raft::device_matrix_view centroids, raft::device_vector_view labels, bool normalize_weight, - raft::host_scalar_view inertia) + raft::host_scalar_view inertia) { cuvs::cluster::kmeans::detail::kmeans_predict( handle, params, X, sample_weight, centroids, labels, normalize_weight, inertia); @@ -227,8 +227,8 @@ void fit_predict(raft::resources const& handle, std::optional> sample_weight, std::optional> centroids, raft::device_vector_view labels, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter) + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter) { cuvs::cluster::kmeans::detail::kmeans_fit_predict( handle, params, X, sample_weight, centroids, labels, inertia, n_iter); @@ -293,9 +293,9 @@ void transform(raft::resources const& handle, * * raft::random::make_blobs(handle, X, labels, n_clusters); * - * auto best_k = raft::make_host_scalar(0); - * auto n_iter = raft::make_host_scalar(0); - * auto inertia = raft::make_host_scalar(0); + * auto best_k = raft::make_host_scalar(0); + * auto n_iter = raft::make_host_scalar(0); + * auto inertia = raft::make_host_scalar(0); * * kmeans::find_k(handle, X, best_k.view(), inertia.view(), n_iter.view(), n_clusters+1); * @@ -316,9 +316,9 @@ void transform(raft::resources const& handle, template void find_k(raft::resources const& handle, raft::device_matrix_view X, - raft::host_scalar_view best_k, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter, + raft::host_scalar_view best_k, + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter, idx_t kmax, idx_t kmin = 1, idx_t maxiter = 100, @@ -378,7 +378,7 @@ template void cluster_cost(raft::resources const& handle, raft::device_vector_view minClusterDistance, rmm::device_uvector& workspace, - raft::device_scalar_view clusterCost, + raft::device_scalar_view clusterCost, ReductionOpT reduction_op) { cuvs::cluster::kmeans::detail::computeClusterCost( @@ -469,7 +469,7 @@ template void cluster_cost(raft::resources const& handle, raft::device_matrix_view X, raft::device_matrix_view centroids, - raft::host_scalar_view cost) + raft::host_scalar_view cost) { auto stream = raft::resource::get_cuda_stream(handle); @@ -681,8 +681,8 @@ void fit_main(raft::resources const& handle, raft::device_matrix_view X, raft::device_vector_view sample_weights, raft::device_matrix_view centroids, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter, + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter, rmm::device_uvector& workspace) { cuvs::cluster::kmeans::detail::kmeans_fit_main( @@ -728,8 +728,8 @@ void kmeans_fit(raft::resources const& handle, raft::device_matrix_view X, std::optional> sample_weight, raft::device_matrix_view centroids, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter) + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter) { kmeans::fit(handle, params, X, sample_weight, centroids, inertia, n_iter); } @@ -777,7 +777,7 @@ void kmeans_predict(raft::resources const& handle, raft::device_matrix_view centroids, raft::device_vector_view labels, bool normalize_weight, - raft::host_scalar_view inertia) + raft::host_scalar_view inertia) { kmeans::predict( handle, params, X, sample_weight, centroids, labels, normalize_weight, inertia); @@ -841,8 +841,8 @@ void kmeans_fit_predict(raft::resources const& handle, std::optional> sample_weight, std::optional> centroids, raft::device_vector_view labels, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter) + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter) { kmeans::fit_predict( handle, params, X, sample_weight, centroids, labels, inertia, n_iter); @@ -1168,8 +1168,8 @@ void kmeans_fit_main(raft::resources const& handle, raft::device_matrix_view X, raft::device_vector_view weight, raft::device_matrix_view centroidsRawData, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter, + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter, rmm::device_uvector& workspace) { kmeans::fit_main( diff --git a/cpp/src/cluster/kmeans_auto_find_k_float.cu b/cpp/src/cluster/kmeans_auto_find_k_float.cu index 44d18f0a11..cb59424422 100644 --- a/cpp/src/cluster/kmeans_auto_find_k_float.cu +++ b/cpp/src/cluster/kmeans_auto_find_k_float.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -23,9 +23,9 @@ namespace cuvs::cluster::kmeans::helpers { void find_k(raft::resources const& handle, raft::device_matrix_view X, - raft::host_scalar_view best_k, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter, + raft::host_scalar_view best_k, + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter, int kmax, int kmin, int maxiter, diff --git a/cpp/src/cluster/kmeans_fit_predict_double.cu b/cpp/src/cluster/kmeans_fit_predict_double.cu index 28a1d70c0c..cfd81ad30e 100644 --- a/cpp/src/cluster/kmeans_fit_predict_double.cu +++ b/cpp/src/cluster/kmeans_fit_predict_double.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,8 +25,8 @@ void fit_predict(raft::resources const& handle, std::optional> sample_weight, std::optional> centroids, raft::device_vector_view labels, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter) + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter) { cuvs::cluster::kmeans::fit_predict( @@ -39,8 +39,8 @@ void fit_predict(raft::resources const& handle, std::optional> sample_weight, std::optional> centroids, raft::device_vector_view labels, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter) + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter) { cuvs::cluster::kmeans::fit_predict( diff --git a/cpp/src/cluster/kmeans_fit_predict_float.cu b/cpp/src/cluster/kmeans_fit_predict_float.cu index be3652db57..6457c26fae 100644 --- a/cpp/src/cluster/kmeans_fit_predict_float.cu +++ b/cpp/src/cluster/kmeans_fit_predict_float.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,8 +25,8 @@ void fit_predict(raft::resources const& handle, std::optional> sample_weight, std::optional> centroids, raft::device_vector_view labels, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter) + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter) { cuvs::cluster::kmeans::fit_predict( @@ -39,8 +39,8 @@ void fit_predict(raft::resources const& handle, std::optional> sample_weight, std::optional> centroids, raft::device_vector_view labels, - raft::host_scalar_view inertia, - raft::host_scalar_view n_iter) + raft::host_scalar_view inertia, + raft::host_scalar_view n_iter) { cuvs::cluster::kmeans::fit_predict( diff --git a/cpp/src/cluster/kmeans_predict_double.cu b/cpp/src/cluster/kmeans_predict_double.cu index 1fcc393acb..deb27af11b 100644 --- a/cpp/src/cluster/kmeans_predict_double.cu +++ b/cpp/src/cluster/kmeans_predict_double.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -26,7 +26,7 @@ void predict(raft::resources const& handle, raft::device_matrix_view centroids, raft::device_vector_view labels, bool normalize_weight, - raft::host_scalar_view inertia) + raft::host_scalar_view inertia) { cuvs::cluster::kmeans::predict( @@ -40,7 +40,7 @@ void predict(raft::resources const& handle, raft::device_matrix_view centroids, raft::device_vector_view labels, bool normalize_weight, - raft::host_scalar_view inertia) + raft::host_scalar_view inertia) { cuvs::cluster::kmeans::predict( diff --git a/cpp/src/cluster/kmeans_predict_float.cu b/cpp/src/cluster/kmeans_predict_float.cu index b5f9f9e510..6c4820f4ff 100644 --- a/cpp/src/cluster/kmeans_predict_float.cu +++ b/cpp/src/cluster/kmeans_predict_float.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -26,7 +26,7 @@ void predict(raft::resources const& handle, raft::device_matrix_view centroids, raft::device_vector_view labels, bool normalize_weight, - raft::host_scalar_view inertia) + raft::host_scalar_view inertia) { cuvs::cluster::kmeans::predict( @@ -39,7 +39,7 @@ void predict(raft::resources const& handle, raft::device_matrix_view centroids, raft::device_vector_view labels, bool normalize_weight, - raft::host_scalar_view inertia) + raft::host_scalar_view inertia) { cuvs::cluster::kmeans::predict( diff --git a/cpp/src/neighbors/cagra.cuh b/cpp/src/neighbors/cagra.cuh index aaf1edb4e3..1e81fe4023 100644 --- a/cpp/src/neighbors/cagra.cuh +++ b/cpp/src/neighbors/cagra.cuh @@ -179,7 +179,7 @@ void build_knn_graph( */ template , + typename accessor = raft::host_device_accessor, raft::memory_type::device>> void build_knn_graph( raft::resources const& res, @@ -222,13 +222,12 @@ void build_knn_graph( * @param[in,out] knn_graph a matrix view (host or device) of the input knn graph [n_rows, * knn_graph_degree] */ -template < - typename DataT, - typename IdxT = uint32_t, - typename d_accessor = raft::host_device_accessor, - raft::memory_type::device>, - typename g_accessor = - raft::host_device_accessor, raft::memory_type::host>> +template , + raft::memory_type::device>, + typename g_accessor = + raft::host_device_accessor, raft::memory_type::host>> void sort_knn_graph( raft::resources const& res, cuvs::distance::DistanceType metric, @@ -238,8 +237,7 @@ void sort_knn_graph( using internal_IdxT = typename std::make_unsigned::type; using g_accessor_internal = - raft::host_device_accessor, - g_accessor::mem_type>; + raft::host_device_accessor, g_accessor::mem_type>; auto knn_graph_internal = raft::mdspan, raft::row_major, g_accessor_internal>( reinterpret_cast(knn_graph.data_handle()), @@ -267,10 +265,9 @@ void sort_knn_graph( * knn_graph_degree] * @param[out] new_graph a host matrix view of the optimized knn graph [n_rows, graph_degree] */ -template < - typename IdxT = uint32_t, - typename g_accessor = - raft::host_device_accessor, raft::memory_type::host>> +template , raft::memory_type::host>> void optimize( raft::resources const& res, raft::mdspan, raft::row_major, g_accessor> knn_graph, @@ -281,9 +278,9 @@ void optimize( } template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> index build( raft::resources const& res, const index_params& params, diff --git a/cpp/src/neighbors/detail/cagra/add_nodes.cuh b/cpp/src/neighbors/detail/cagra/add_nodes.cuh index 8023d8f9ef..1f52b7fa69 100644 --- a/cpp/src/neighbors/detail/cagra/add_nodes.cuh +++ b/cpp/src/neighbors/detail/cagra/add_nodes.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -471,7 +471,7 @@ void extend_core( using out_owning_type = owning_dataset; auto out_layout = raft::make_strided_layout(updated_dataset_view.extents(), - std::array{stride, 1}); + cuda::std::array{stride, 1}); index.update_dataset(handle, out_owning_type{std::move(updated_dataset), out_layout}); } diff --git a/cpp/src/neighbors/detail/cagra/cagra_build.cuh b/cpp/src/neighbors/detail/cagra/cagra_build.cuh index e2a43b8a07..9a2a7ca607 100644 --- a/cpp/src/neighbors/detail/cagra/cagra_build.cuh +++ b/cpp/src/neighbors/detail/cagra/cagra_build.cuh @@ -323,7 +323,7 @@ void build_knn_graph( gpu_top_k); } } else { - auto neighbor_candidates_view = raft::make_device_matrix_view( + auto neighbor_candidates_view = raft::make_device_matrix_view( neighbors.data_handle(), batch.size(), gpu_top_k); auto refined_neighbors_view = raft::make_device_matrix_view( refined_neighbors.data_handle(), batch.size(), top_k); @@ -397,8 +397,7 @@ void build_knn_graph( using internal_IdxT = typename std::make_unsigned::type; using g_accessor = typename decltype(nn_descent_idx.graph())::accessor_type; using g_accessor_internal = - raft::host_device_accessor, - g_accessor::mem_type>; + raft::host_device_accessor, g_accessor::mem_type>; auto knn_graph_internal = raft::mdspan, raft::row_major, g_accessor_internal>( @@ -410,10 +409,9 @@ void build_knn_graph( res, build_params.metric, dataset, knn_graph_internal); } -template < - typename IdxT = uint32_t, - typename g_accessor = - raft::host_device_accessor, raft::memory_type::host>> +template , raft::memory_type::host>> void optimize( raft::resources const& res, raft::mdspan, raft::row_major, g_accessor> knn_graph, @@ -428,8 +426,7 @@ void optimize( new_graph.extent(1)); using g_accessor_internal = - raft::host_device_accessor, - raft::memory_type::host>; + raft::host_device_accessor, raft::memory_type::host>; auto knn_graph_internal = raft::mdspan, raft::row_major, g_accessor_internal>( reinterpret_cast(knn_graph.data_handle()), @@ -487,9 +484,9 @@ struct mmap_owner { }; template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> auto iterative_build_graph( raft::resources const& res, const index_params& params, @@ -658,9 +655,9 @@ auto iterative_build_graph( } template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> index build( raft::resources const& res, const index_params& params, diff --git a/cpp/src/neighbors/detail/cagra/cagra_merge.cuh b/cpp/src/neighbors/detail/cagra/cagra_merge.cuh index 86faad02c0..41a84a068c 100644 --- a/cpp/src/neighbors/detail/cagra/cagra_merge.cuh +++ b/cpp/src/neighbors/detail/cagra/cagra_merge.cuh @@ -111,7 +111,7 @@ index merge(raft::resources const& handle, using container_policy_t = typename matrix_t::container_policy_type; using owning_t = owning_dataset; auto out_layout = raft::make_strided_layout(updated_dataset.view().extents(), - std::array{stride, 1}); + cuda::std::array{stride, 1}); merged_index.update_dataset(handle, owning_t{std::move(updated_dataset), out_layout}); } RAFT_LOG_DEBUG("cagra merge: using device memory for merged dataset"); @@ -133,7 +133,7 @@ index merge(raft::resources const& handle, using container_policy_t = typename matrix_t::container_policy_type; using owning_t = owning_dataset; auto out_layout = raft::make_strided_layout(updated_dataset.view().extents(), - std::array{stride, 1}); + cuda::std::array{stride, 1}); merged_index.update_dataset(handle, owning_t{std::move(updated_dataset), out_layout}); } return merged_index; diff --git a/cpp/src/neighbors/detail/cagra/graph_core.cuh b/cpp/src/neighbors/detail/cagra/graph_core.cuh index 809e3b6dd4..3ea2a52e91 100644 --- a/cpp/src/neighbors/detail/cagra/graph_core.cuh +++ b/cpp/src/neighbors/detail/cagra/graph_core.cuh @@ -499,13 +499,12 @@ void shift_array(T* array, uint64_t num) } } // namespace -template < - typename DataT, - typename IdxT = uint32_t, - typename d_accessor = raft::host_device_accessor, - raft::memory_type::device>, - typename g_accessor = - raft::host_device_accessor, raft::memory_type::host>> +template , + raft::memory_type::device>, + typename g_accessor = + raft::host_device_accessor, raft::memory_type::host>> void sort_knn_graph( raft::resources const& res, const cuvs::distance::DistanceType metric, @@ -1167,10 +1166,9 @@ void count_2hop_detours(raft::host_matrix_view k } } -template < - typename IdxT = uint32_t, - typename g_accessor = - raft::host_device_accessor, raft::memory_type::host>> +template , raft::memory_type::host>> void optimize( raft::resources const& res, raft::mdspan, raft::row_major, g_accessor> knn_graph, diff --git a/cpp/src/neighbors/detail/knn_brute_force.cuh b/cpp/src/neighbors/detail/knn_brute_force.cuh index 7da7a0916d..c048c69bbe 100644 --- a/cpp/src/neighbors/detail/knn_brute_force.cuh +++ b/cpp/src/neighbors/detail/knn_brute_force.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -530,7 +530,7 @@ void brute_force_knn_impl( raft::make_device_matrix_view(out_I, n, input.size() * k), raft::make_device_matrix_view(res_D, n, k), raft::make_device_matrix_view(res_I, n, k), - raft::make_device_vector_view(trans.data(), input.size())); + raft::make_device_vector_view(trans.data(), static_cast(input.size()))); } }; diff --git a/cpp/src/neighbors/detail/nn_descent.cuh b/cpp/src/neighbors/detail/nn_descent.cuh index 776f5c417d..f298d33d78 100644 --- a/cpp/src/neighbors/detail/nn_descent.cuh +++ b/cpp/src/neighbors/detail/nn_descent.cuh @@ -1309,9 +1309,9 @@ void GNND::build(Data_t* data, } template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> void build(raft::resources const& res, const index_params& params, raft::mdspan, raft::row_major, Accessor> dataset, @@ -1355,9 +1355,9 @@ void build(raft::resources const& res, } template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> index build( raft::resources const& res, const index_params& params, diff --git a/cpp/src/neighbors/detail/vamana/greedy_search.cuh b/cpp/src/neighbors/detail/vamana/greedy_search.cuh index 9e884fc3ea..31670f3d13 100644 --- a/cpp/src/neighbors/detail/vamana/greedy_search.cuh +++ b/cpp/src/neighbors/detail/vamana/greedy_search.cuh @@ -93,9 +93,9 @@ __global__ void SortPairsKernel(void* query_list_ptr, int num_queries, int topk) **********************************************************************************************/ template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> __global__ void GreedySearchKernel( raft::device_matrix_view graph, raft::mdspan, raft::row_major, Accessor> dataset, diff --git a/cpp/src/neighbors/detail/vamana/robust_prune.cuh b/cpp/src/neighbors/detail/vamana/robust_prune.cuh index 0a76093444..e37e9474f9 100644 --- a/cpp/src/neighbors/detail/vamana/robust_prune.cuh +++ b/cpp/src/neighbors/detail/vamana/robust_prune.cuh @@ -62,9 +62,9 @@ namespace { **********************************************************************************************/ template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> __global__ void RobustPruneKernel( raft::device_matrix_view graph, raft::mdspan, raft::row_major, Accessor> dataset, diff --git a/cpp/src/neighbors/detail/vamana/vamana_build.cuh b/cpp/src/neighbors/detail/vamana/vamana_build.cuh index 629319e0e3..d6d6c765d0 100644 --- a/cpp/src/neighbors/detail/vamana/vamana_build.cuh +++ b/cpp/src/neighbors/detail/vamana/vamana_build.cuh @@ -105,9 +105,9 @@ __global__ void print_queryIds(void* query_list_ptr) *******************************************************************************************/ template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> void batched_insert_vamana( raft::resources const& res, const index_params& params, @@ -578,9 +578,9 @@ auto quantize_all_vectors(raft::resources const& res, } template , - raft::memory_type::host>> + typename IdxT = uint64_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> index build( raft::resources const& res, const index_params& params, diff --git a/cpp/src/neighbors/detail/vamana/vamana_structs.cuh b/cpp/src/neighbors/detail/vamana/vamana_structs.cuh index d5812e1ff6..b6344641d6 100644 --- a/cpp/src/neighbors/detail/vamana/vamana_structs.cuh +++ b/cpp/src/neighbors/detail/vamana/vamana_structs.cuh @@ -451,9 +451,9 @@ __global__ void populate_reverse_list_struct(QueryCandidates* revers // Recompute distances of reverse list. Allows us to avoid keeping distances during sort template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> __global__ void recompute_reverse_dists( QueryCandidates* reverse_list, raft::mdspan, raft::row_major, Accessor> dataset, diff --git a/cpp/src/neighbors/iface/generate_iface.py b/cpp/src/neighbors/iface/generate_iface.py index 66209e3f47..e5dd34c6b6 100644 --- a/cpp/src/neighbors/iface/generate_iface.py +++ b/cpp/src/neighbors/iface/generate_iface.py @@ -53,10 +53,10 @@ flat_macro = """ #define CUVS_INST_MG_FLAT(T, IdxT) \\ - using T_ha = raft::host_device_accessor, raft::memory_type::device>; \\ - using T_da= raft::host_device_accessor, raft::memory_type::host>; \\ - using IdxT_ha = raft::host_device_accessor, raft::memory_type::device>; \\ - using IdxT_da = raft::host_device_accessor, raft::memory_type::host>; \\ + using T_ha = raft::host_device_accessor, raft::memory_type::device>; \\ + using T_da = raft::host_device_accessor, raft::memory_type::host>; \\ + using IdxT_ha = raft::host_device_accessor, raft::memory_type::device>; \\ + using IdxT_da = raft::host_device_accessor, raft::memory_type::host>; \\ \\ template void build(const raft::resources& handle, \\ cuvs::neighbors::iface, T, IdxT>& interface, \\ @@ -107,10 +107,10 @@ pq_macro = """ #define CUVS_INST_MG_PQ(T, IdxT) \\ - using T_ha = raft::host_device_accessor, raft::memory_type::device>; \\ - using T_da= raft::host_device_accessor, raft::memory_type::host>; \\ - using IdxT_ha = raft::host_device_accessor, raft::memory_type::device>; \\ - using IdxT_da = raft::host_device_accessor, raft::memory_type::host>; \\ + using T_ha = raft::host_device_accessor, raft::memory_type::device>; \\ + using T_da = raft::host_device_accessor, raft::memory_type::host>; \\ + using IdxT_ha = raft::host_device_accessor, raft::memory_type::device>; \\ + using IdxT_da = raft::host_device_accessor, raft::memory_type::host>; \\ \\ template void build(const raft::resources& handle, \\ cuvs::neighbors::iface, T, IdxT>& interface, \\ @@ -161,10 +161,10 @@ cagra_macro = """ #define CUVS_INST_MG_CAGRA(T, IdxT) \\ - using T_ha = raft::host_device_accessor, raft::memory_type::device>; \\ - using T_da= raft::host_device_accessor, raft::memory_type::host>; \\ - using IdxT_ha = raft::host_device_accessor, raft::memory_type::device>; \\ - using IdxT_da = raft::host_device_accessor, raft::memory_type::host>; \\ + using T_ha = raft::host_device_accessor, raft::memory_type::device>; \\ + using T_da = raft::host_device_accessor, raft::memory_type::host>; \\ + using IdxT_ha = raft::host_device_accessor, raft::memory_type::device>; \\ + using IdxT_da = raft::host_device_accessor, raft::memory_type::host>; \\ \\ template void build(const raft::resources& handle, \\ cuvs::neighbors::iface, T, IdxT>& interface, \\ diff --git a/cpp/src/neighbors/iface/iface_cagra_float_uint32_t.cu b/cpp/src/neighbors/iface/iface_cagra_float_uint32_t.cu index dae6508558..8f1bab435a 100644 --- a/cpp/src/neighbors/iface/iface_cagra_float_uint32_t.cu +++ b/cpp/src/neighbors/iface/iface_cagra_float_uint32_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_CAGRA(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_cagra_half_uint32_t.cu b/cpp/src/neighbors/iface/iface_cagra_half_uint32_t.cu index 7219dd3ad9..52e68ffff0 100644 --- a/cpp/src/neighbors/iface/iface_cagra_half_uint32_t.cu +++ b/cpp/src/neighbors/iface/iface_cagra_half_uint32_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_CAGRA(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_cagra_int8_t_uint32_t.cu b/cpp/src/neighbors/iface/iface_cagra_int8_t_uint32_t.cu index 227b03eb11..2f5d3116f9 100644 --- a/cpp/src/neighbors/iface/iface_cagra_int8_t_uint32_t.cu +++ b/cpp/src/neighbors/iface/iface_cagra_int8_t_uint32_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_CAGRA(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_cagra_uint8_t_uint32_t.cu b/cpp/src/neighbors/iface/iface_cagra_uint8_t_uint32_t.cu index b55d96c419..931ecfafac 100644 --- a/cpp/src/neighbors/iface/iface_cagra_uint8_t_uint32_t.cu +++ b/cpp/src/neighbors/iface/iface_cagra_uint8_t_uint32_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_CAGRA(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_flat_float_int64_t.cu b/cpp/src/neighbors/iface/iface_flat_float_int64_t.cu index 9d7e40db15..baff52fd3a 100644 --- a/cpp/src/neighbors/iface/iface_flat_float_int64_t.cu +++ b/cpp/src/neighbors/iface/iface_flat_float_int64_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_FLAT(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_flat_half_int64_t.cu b/cpp/src/neighbors/iface/iface_flat_half_int64_t.cu index 4907e11063..3d4acb19c8 100644 --- a/cpp/src/neighbors/iface/iface_flat_half_int64_t.cu +++ b/cpp/src/neighbors/iface/iface_flat_half_int64_t.cu @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_FLAT(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_flat_int8_t_int64_t.cu b/cpp/src/neighbors/iface/iface_flat_int8_t_int64_t.cu index 8daa6b435d..e435a1daf2 100644 --- a/cpp/src/neighbors/iface/iface_flat_int8_t_int64_t.cu +++ b/cpp/src/neighbors/iface/iface_flat_int8_t_int64_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_FLAT(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_flat_uint8_t_int64_t.cu b/cpp/src/neighbors/iface/iface_flat_uint8_t_int64_t.cu index 0fdaef195d..d1061a5177 100644 --- a/cpp/src/neighbors/iface/iface_flat_uint8_t_int64_t.cu +++ b/cpp/src/neighbors/iface/iface_flat_uint8_t_int64_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_FLAT(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_pq_float_int64_t.cu b/cpp/src/neighbors/iface/iface_pq_float_int64_t.cu index c36d92d859..aed529230b 100644 --- a/cpp/src/neighbors/iface/iface_pq_float_int64_t.cu +++ b/cpp/src/neighbors/iface/iface_pq_float_int64_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_PQ(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_pq_half_int64_t.cu b/cpp/src/neighbors/iface/iface_pq_half_int64_t.cu index 3f8ec9ac42..f5fef82f4f 100644 --- a/cpp/src/neighbors/iface/iface_pq_half_int64_t.cu +++ b/cpp/src/neighbors/iface/iface_pq_half_int64_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_PQ(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_pq_int8_t_int64_t.cu b/cpp/src/neighbors/iface/iface_pq_int8_t_int64_t.cu index e58132b5b0..d2223238a6 100644 --- a/cpp/src/neighbors/iface/iface_pq_int8_t_int64_t.cu +++ b/cpp/src/neighbors/iface/iface_pq_int8_t_int64_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_PQ(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/iface/iface_pq_uint8_t_int64_t.cu b/cpp/src/neighbors/iface/iface_pq_uint8_t_int64_t.cu index b6ce641b54..4d9032fd78 100644 --- a/cpp/src/neighbors/iface/iface_pq_uint8_t_int64_t.cu +++ b/cpp/src/neighbors/iface/iface_pq_uint8_t_int64_t.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,14 +28,14 @@ namespace cuvs::neighbors { #define CUVS_INST_MG_PQ(T, IdxT) \ - using T_ha = raft::host_device_accessor, \ + using T_ha = \ + raft::host_device_accessor, raft::memory_type::device>; \ + using T_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ + using IdxT_ha = raft::host_device_accessor, \ raft::memory_type::device>; \ - using T_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ - using IdxT_ha = raft::host_device_accessor, \ - raft::memory_type::device>; \ - using IdxT_da = raft::host_device_accessor, \ - raft::memory_type::host>; \ + using IdxT_da = \ + raft::host_device_accessor, raft::memory_type::host>; \ \ template void build( \ const raft::resources& handle, \ diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_build_common.cu b/cpp/src/neighbors/ivf_pq/ivf_pq_build_common.cu index dc54964275..8178614170 100644 --- a/cpp/src/neighbors/ivf_pq/ivf_pq_build_common.cu +++ b/cpp/src/neighbors/ivf_pq/ivf_pq_build_common.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2024, NVIDIA CORPORATION. + * Copyright (c) 2022-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -280,7 +280,7 @@ void make_rotation_matrix(raft::resources const& res, void set_centers(raft::resources const& handle, index* index, - raft::device_matrix_view cluster_centers) + raft::device_matrix_view cluster_centers) { RAFT_EXPECTS(cluster_centers.extent(0) == index->n_lists(), "Number of rows in the new centers must be equal to the number of IVF lists"); @@ -292,14 +292,14 @@ void set_centers(raft::resources const& handle, void extract_centers(raft::resources const& res, const cuvs::neighbors::ivf_pq::index& index, - raft::device_matrix_view cluster_centers) + raft::device_matrix_view cluster_centers) { detail::extract_centers(res, index, cluster_centers); } void extract_centers(raft::resources const& res, const cuvs::neighbors::ivf_pq::index& index, - raft::host_matrix_view cluster_centers) + raft::host_matrix_view cluster_centers) { detail::extract_centers(res, index, cluster_centers); } diff --git a/cpp/src/neighbors/scann/detail/scann_avq.cuh b/cpp/src/neighbors/scann/detail/scann_avq.cuh index 1c42f6d2b9..d7a439cf36 100644 --- a/cpp/src/neighbors/scann/detail/scann_avq.cuh +++ b/cpp/src/neighbors/scann/detail/scann_avq.cuh @@ -236,7 +236,7 @@ void compute_avq_centroid(raft::resources const& dev_resources, // Compute || x_i || ^ 0.5 * (eta -3) auto norms_eta_3 = raft::make_device_vector(dev_resources, x.extent(0)); - auto eta_3 = raft::make_host_scalar((eta - 3) / 2); + auto eta_3 = raft::make_host_scalar((eta - 3) / 2); raft::linalg::power_scalar(dev_resources, raft::make_const_mdspan(norms.view()), @@ -245,14 +245,14 @@ void compute_avq_centroid(raft::resources const& dev_resources, // Compute || x_i || ^ (eta - 1) auto norms_eta_1 = raft::make_device_vector(dev_resources, x.extent(0)); - auto eta_1 = raft::make_host_scalar(eta - 1); + auto eta_1 = raft::make_host_scalar(eta - 1); raft::linalg::power_scalar(dev_resources, raft::make_const_mdspan(norms.view()), norms_eta_1.view(), raft::make_const_mdspan(eta_1.view())); - auto sum_norms_eta_1 = raft::make_device_scalar(dev_resources, 0.0); + auto sum_norms_eta_1 = raft::make_device_scalar(dev_resources, 0.0); sum_reduce_vector(dev_resources, norms_eta_1.view(), sum_norms_eta_1.view()); @@ -294,7 +294,7 @@ void compute_avq_centroid(raft::resources const& dev_resources, raft::matrix::eye(dev_resources, x_trans_x.view()); - auto eta_m_1 = raft::make_device_scalar(dev_resources, eta - 1); + auto eta_m_1 = raft::make_device_scalar(dev_resources, eta - 1); auto cublas_handle = raft::resource::get_cublas_handle(dev_resources); raft::linalg::detail::cublas_device_pointer_mode pm(cublas_handle); @@ -318,7 +318,7 @@ void compute_avq_centroid(raft::resources const& dev_resources, cholesky_solver(dev_resources, x_trans_x.view(), avq_centroid, avq_centroid); - auto h_eta = raft::make_host_scalar(eta); + auto h_eta = raft::make_host_scalar(eta); raft::linalg::multiply_scalar(dev_resources, raft::make_const_mdspan(avq_centroid), @@ -357,8 +357,8 @@ void rescale_avq_centroids(raft::resources const& dev_resources, raft::device_vector_view cluster_sizes, uint32_t dataset_size) { - auto rescale_num = raft::make_device_scalar(dev_resources, 0); - auto rescale_denom = raft::make_device_scalar(dev_resources, 0); + auto rescale_num = raft::make_device_scalar(dev_resources, 0); + auto rescale_denom = raft::make_device_scalar(dev_resources, 0); sum_reduce_vector(dev_resources, rescale_num_v, rescale_num.view()); @@ -411,8 +411,8 @@ class cluster_loader { raft::device_matrix d_cluster_buf_; raft::device_matrix d_cluster_copy_buf_; const T* dataset_ptr_; - raft::host_vector_view h_cluster_offsets_; - raft::device_vector_view cluster_ids_; + raft::host_vector_view h_cluster_offsets_; + raft::device_vector_view cluster_ids_; cudaStream_t stream_; int64_t dim_; int64_t n_rows_; @@ -435,8 +435,8 @@ class cluster_loader { int64_t n_rows, int64_t max_cluster_size, int64_t h_buf_size, - raft::host_vector_view h_cluster_offsets, - raft::device_vector_view cluster_ids, + raft::host_vector_view h_cluster_offsets, + raft::device_vector_view cluster_ids, bool needs_copy, cudaStream_t stream) : dim_(dim), @@ -456,8 +456,8 @@ class cluster_loader { public: cluster_loader(raft::resources const& res, raft::device_matrix_view dataset_view, - raft::host_vector_view h_cluster_offsets, - raft::device_vector_view cluster_ids, + raft::host_vector_view h_cluster_offsets, + raft::device_vector_view cluster_ids, int64_t max_cluster_size, cudaStream_t stream) : cluster_loader(res, @@ -476,8 +476,8 @@ class cluster_loader { cluster_loader(raft::resources const& res, raft::host_matrix_view dataset_view, - raft::host_vector_view h_cluster_offsets, - raft::device_vector_view cluster_ids, + raft::host_vector_view h_cluster_offsets, + raft::device_vector_view cluster_ids, int64_t max_cluster_size, cudaStream_t stream) : cluster_loader(res, @@ -593,10 +593,10 @@ class cluster_loader { * @param eta the weight for the parallel component of the residual in the avq update */ template , - raft::memory_type::host>> + typename IdxT = int64_t, + typename LabelT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> void apply_avq(raft::resources const& res, raft::mdspan, raft::row_major, Accessor> dataset, raft::device_matrix_view centroids_view, diff --git a/cpp/src/neighbors/scann/detail/scann_build.cuh b/cpp/src/neighbors/scann/detail/scann_build.cuh index 62ab8a317f..1fd02345c3 100644 --- a/cpp/src/neighbors/scann/detail/scann_build.cuh +++ b/cpp/src/neighbors/scann/detail/scann_build.cuh @@ -85,9 +85,9 @@ using namespace cuvs::spatial::knn::detail; // NOLINT static const std::string RAFT_NAME = "raft"; template , - raft::memory_type::host>> + typename IdxT = int64_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> index build( raft::resources const& res, const index_params& params, diff --git a/cpp/src/neighbors/scann/detail/scann_quantize.cuh b/cpp/src/neighbors/scann/detail/scann_quantize.cuh index 3cc0f40ebf..8fd168733f 100644 --- a/cpp/src/neighbors/scann/detail/scann_quantize.cuh +++ b/cpp/src/neighbors/scann/detail/scann_quantize.cuh @@ -283,8 +283,8 @@ void unpack_codes(raft::resources const& res, */ template , - raft::memory_type::host>> + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> auto sample_training_residuals( raft::resources const& res, random::RngState random_state, diff --git a/cpp/src/neighbors/scann/detail/scann_serialize.cuh b/cpp/src/neighbors/scann/detail/scann_serialize.cuh index 4a8788d679..08457c8df1 100644 --- a/cpp/src/neighbors/scann/detail/scann_serialize.cuh +++ b/cpp/src/neighbors/scann/detail/scann_serialize.cuh @@ -50,8 +50,8 @@ std::ofstream open_file(std::string file_name) // Helper for serializing device/host matrix to a given file template , - raft::memory_type::host>> + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> void serialize_matrix( raft::resources const& res, std::filesystem::path file_path, diff --git a/cpp/src/neighbors/scann/scann.cuh b/cpp/src/neighbors/scann/scann.cuh index d45790aa75..5825af0d69 100644 --- a/cpp/src/neighbors/scann/scann.cuh +++ b/cpp/src/neighbors/scann/scann.cuh @@ -64,9 +64,9 @@ namespace cuvs::neighbors::experimental::scann { * @return the constructed scann index */ template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> index build( raft::resources const& res, const index_params& params, diff --git a/cpp/src/neighbors/vamana.cuh b/cpp/src/neighbors/vamana.cuh index af8cc432d1..0742fe755b 100644 --- a/cpp/src/neighbors/vamana.cuh +++ b/cpp/src/neighbors/vamana.cuh @@ -78,9 +78,9 @@ namespace cuvs::neighbors::vamana { * @return the constructed vamana index */ template , - raft::memory_type::host>> + typename IdxT = uint32_t, + typename Accessor = + raft::host_device_accessor, raft::memory_type::host>> index build( raft::resources const& res, const index_params& params, diff --git a/cpp/src/sparse/cluster/cluster_solvers.cuh b/cpp/src/sparse/cluster/cluster_solvers.cuh index 737dfa5f8a..9b319b2cab 100644 --- a/cpp/src/sparse/cluster/cluster_solvers.cuh +++ b/cpp/src/sparse/cluster/cluster_solvers.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2024, NVIDIA CORPORATION. + * Copyright (c) 2019-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -66,8 +66,8 @@ struct kmeans_solver_t { km_params.max_iter = config_.maxIter; km_params.rng_state.seed = config_.seed; - auto X = raft::make_device_matrix_view(obs, n_obs_vecs, dim); - auto labels = raft::make_device_vector_view(codes, n_obs_vecs); + auto X = raft::make_device_matrix_view(obs, n_obs_vecs, dim); + auto labels = raft::make_device_vector_view(codes, n_obs_vecs); auto centroids = raft::make_device_matrix(handle, config_.n_clusters, dim); auto weight = raft::make_device_vector(handle, n_obs_vecs); @@ -76,15 +76,17 @@ struct kmeans_solver_t { weight.data_handle() + n_obs_vecs, 1); - auto sw = std::make_optional((raft::device_vector_view)weight.view()); - cuvs::cluster::kmeans::fit_predict(handle, - km_params, - X, - sw, - centroids.view(), - labels, - raft::make_host_scalar_view(&residual), - raft::make_host_scalar_view(&iters)); + auto sw = + std::make_optional((raft::device_vector_view)weight.view()); + cuvs::cluster::kmeans::fit_predict( + handle, + km_params, + X, + sw, + centroids.view(), + labels, + raft::make_host_scalar_view(&residual), + raft::make_host_scalar_view(&iters)); return std::make_pair(residual, iters); } diff --git a/cpp/tests/cluster/kmeans.cu b/cpp/tests/cluster/kmeans.cu index 7b6cfda28b..3f170219ce 100644 --- a/cpp/tests/cluster/kmeans.cu +++ b/cpp/tests/cluster/kmeans.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2024, NVIDIA CORPORATION. + * Copyright (c) 2022-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -164,8 +164,8 @@ class KmeansTest : public ::testing::TestWithParam> { // X_view, // weight_view, // centroids_view, - // raft::make_host_scalar_view(&inertia), - // raft::make_host_scalar_view(&n_iter), + // raft::make_host_scalar_view(&inertia), + // raft::make_host_scalar_view(&n_iter), // workspace); // // // Check that the cluster cost decreased @@ -293,8 +293,8 @@ class KmeansTest : public ::testing::TestWithParam> { d_sw, d_centroids_view, raft::make_device_vector_view(d_labels.data(), n_samples), - raft::make_host_scalar_view(&inertia), - raft::make_host_scalar_view(&n_iter)); + raft::make_host_scalar_view(&inertia), + raft::make_host_scalar_view(&n_iter)); raft::resource::sync_stream(handle, stream); diff --git a/cpp/tests/cluster/kmeans_mg.cu b/cpp/tests/cluster/kmeans_mg.cu index b9e06b2f13..664ca1d9b4 100644 --- a/cpp/tests/cluster/kmeans_mg.cu +++ b/cpp/tests/cluster/kmeans_mg.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2024, NVIDIA CORPORATION. + * Copyright (c) 2022-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -129,8 +129,8 @@ class KmeansTest : public ::testing::TestWithParam> { X_view, d_sw, centroids_view, - raft::make_host_scalar_view(&inertia), - raft::make_host_scalar_view(&n_iter)); + raft::make_host_scalar_view(&inertia), + raft::make_host_scalar_view(&n_iter)); cuvs::cluster::kmeans::predict( handle, @@ -140,7 +140,7 @@ class KmeansTest : public ::testing::TestWithParam> { d_centroids.data(), raft::make_device_vector_view(d_labels.data(), n_samples), true, - raft::make_host_scalar_view(&inertia)); + raft::make_host_scalar_view(&inertia)); score = raft::stats::adjusted_rand_index( d_labels_ref.data(), d_labels.data(), n_samples, raft::resource::get_cuda_stream(handle)); handle.sync_stream(stream);