diff --git a/cpp/src/cluster/detail/kmeans_balanced.cuh b/cpp/src/cluster/detail/kmeans_balanced.cuh index ac0430b430..272d45f13d 100644 --- a/cpp/src/cluster/detail/kmeans_balanced.cuh +++ b/cpp/src/cluster/detail/kmeans_balanced.cuh @@ -681,10 +681,12 @@ auto adjust_centers(const raft::resources& handle, rmm::device_uvector donor_clusters(n_pairs, stream, device_memory); constexpr uint32_t kBlockDimY = 4; const dim3 block_dim(raft::WarpSize, kBlockDimY, 1); - rmm::device_scalar update_count(0, stream, device_memory); + rmm::device_scalar update_count(stream, device_memory); + update_count.set_value_to_zero_async(stream); if (donor_selection == cuvs::cluster::kmeans::balanced_donor_selection::Random) { - rmm::device_scalar search_count(0, stream, device_memory); + rmm::device_scalar search_count(stream, device_memory); + search_count.set_value_to_zero_async(stream); const dim3 grid_dim(raft::ceildiv(n_clusters, static_cast(kBlockDimY)), 1, 1); adjust_centers_random_donor_kernel <<>>(centers, diff --git a/cpp/src/core/bloom_filter.cu b/cpp/src/core/bloom_filter.cu index f8f12d28e3..b25192042c 100644 --- a/cpp/src/core/bloom_filter.cu +++ b/cpp/src/core/bloom_filter.cu @@ -7,6 +7,7 @@ #include #include +#include #include #include @@ -20,7 +21,7 @@ namespace cuvs::core { namespace { -using default_filter_policy = cuco::default_filter_policy; +using default_filter_policy = cuco::bloom_filter_policy; constexpr auto kPatternBits = default_filter_policy::pattern_bits; constexpr auto kWordsPerBlock = default_filter_policy::words_per_block; diff --git a/cpp/src/distance/detail/sparse/coo_spmv_strategies/coo_mask_row_iterators.cuh b/cpp/src/distance/detail/sparse/coo_spmv_strategies/coo_mask_row_iterators.cuh index f518823f64..8f56320a4c 100644 --- a/cpp/src/distance/detail/sparse/coo_spmv_strategies/coo_mask_row_iterators.cuh +++ b/cpp/src/distance/detail/sparse/coo_spmv_strategies/coo_mask_row_iterators.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -118,8 +118,7 @@ class chunked_mask_row_it : public mask_row_it { { auto policy = rmm::exec_policy(stream); - constexpr value_idx first_element = 0; - n_chunks_per_row.set_element_async(0, first_element, stream); + n_chunks_per_row.set_element_to_zero_async(0, stream); n_chunks_per_row_functor chunk_functor(indptr, row_chunk_size); thrust::transform( policy, mask_row_idx, mask_row_idx + n_rows, n_chunks_per_row.begin() + 1, chunk_functor); diff --git a/cpp/tests/neighbors/brute_force_prefiltered.cu b/cpp/tests/neighbors/brute_force_prefiltered.cu index 1c9f9966a7..33ced34693 100644 --- a/cpp/tests/neighbors/brute_force_prefiltered.cu +++ b/cpp/tests/neighbors/brute_force_prefiltered.cu @@ -198,10 +198,11 @@ class PrefilteredBruteForceOnBitmapTest index_t nnz_h = 0; { - auto src = out_src.data(); - auto dst = out_dst.data(); - auto bitmap = filter_d.data(); - rmm::device_scalar nnz(0, stream); + auto src = out_src.data(); + auto dst = out_dst.data(); + auto bitmap = filter_d.data(); + index_t zero = 0; + rmm::device_scalar nnz(zero, stream); auto nnz_view = raft::make_device_scalar_view(nnz.data()); auto filter_view = raft::make_device_vector_view(filter_d.data(), filter_d.size()); @@ -616,10 +617,11 @@ class PrefilteredBruteForceOnBitsetTest index_t nnz_h = 0; { - auto src = out_src.data(); - auto dst = out_dst.data(); - auto bitset = filter_d.data(); - rmm::device_scalar nnz(0, stream); + auto src = out_src.data(); + auto dst = out_dst.data(); + auto bitset = filter_d.data(); + index_t zero = 0; + rmm::device_scalar nnz(zero, stream); auto nnz_view = raft::make_device_scalar_view(nnz.data()); auto filter_view = raft::make_device_vector_view(filter_d.data(), filter_d.size());