Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 4 additions & 2 deletions cpp/src/cluster/detail/kmeans_balanced.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -681,10 +681,12 @@ auto adjust_centers(const raft::resources& handle,
rmm::device_uvector<IdxT> donor_clusters(n_pairs, stream, device_memory);
constexpr uint32_t kBlockDimY = 4;
const dim3 block_dim(raft::WarpSize, kBlockDimY, 1);
rmm::device_scalar<IdxT> update_count(0, stream, device_memory);
rmm::device_scalar<IdxT> 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<IdxT> search_count(0, stream, device_memory);
rmm::device_scalar<IdxT> search_count(stream, device_memory);
search_count.set_value_to_zero_async(stream);
const dim3 grid_dim(raft::ceildiv(n_clusters, static_cast<IdxT>(kBlockDimY)), 1, 1);
adjust_centers_random_donor_kernel<kBlockDimY>
<<<grid_dim, block_dim, 0, stream>>>(centers,
Expand Down
3 changes: 2 additions & 1 deletion cpp/src/core/bloom_filter.cu
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,7 @@
#include <cuvs/core/bloom_filter.hpp>

#include <cuco/bloom_filter.cuh>
#include <cuco/bloom_filter_policy.cuh>

#include <raft/core/error.hpp>
#include <raft/core/resource/cuda_stream.hpp>
Expand All @@ -20,7 +21,7 @@ namespace cuvs::core {

namespace {

using default_filter_policy = cuco::default_filter_policy<bloom_filter::key_type>;
using default_filter_policy = cuco::bloom_filter_policy<bloom_filter::key_type>;

constexpr auto kPatternBits = default_filter_policy::pattern_bits;
constexpr auto kWordsPerBlock = default_filter_policy::words_per_block;
Expand Down
Original file line number Diff line number Diff line change
@@ -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
*/

Expand Down Expand Up @@ -118,8 +118,7 @@ class chunked_mask_row_it : public mask_row_it<value_idx> {
{
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);
Expand Down
18 changes: 10 additions & 8 deletions cpp/tests/neighbors/brute_force_prefiltered.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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<index_t> 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<index_t> nnz(zero, stream);
auto nnz_view = raft::make_device_scalar_view<index_t>(nnz.data());
auto filter_view =
raft::make_device_vector_view<const uint32_t, index_t>(filter_d.data(), filter_d.size());
Expand Down Expand Up @@ -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<index_t> 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<index_t> nnz(zero, stream);
auto nnz_view = raft::make_device_scalar_view<index_t>(nnz.data());
auto filter_view =
raft::make_device_vector_view<const uint32_t, index_t>(filter_d.data(), filter_d.size());
Expand Down
Loading