Skip to content

IVF-Flat: fix irrelevant assert in the fused kernel mode #1941

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
rapids-bot merged 4 commits into from
Mar 24, 2026
Merged

IVF-Flat: fix irrelevant assert in the fused kernel mode #1941

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
98c88e3
IVF-Flat: fix irrelevant assert in the fused kernel mode
achirkin Mar 6, 2026
b126e2f
Merge branch 'main' into fix-ivf-flat-irrelevant-assert
achirkin Mar 9, 2026
6f3716d
Update ivf_flat_interleaved_scan_kernel.cuh
achirkin Mar 9, 2026
0ea9300
Merge branch 'release/26.04' into fix-ivf-flat-irrelevant-assert
achirkin Mar 24, 2026
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
7 changes: 5 additions & 2 deletions cpp/src/neighbors/ivf_flat/ivf_flat_interleaved_scan.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
* SPDX-License-Identifier: Apache-2.0
*/

Expand Down Expand Up @@ -875,7 +875,10 @@ RAFT_KERNEL __launch_bounds__(kThreadsPerBlock)
uint32_t sample_offset = 0;
if (probe_id > 0) { sample_offset = chunk_indices[probe_id - 1]; }
assert(list_length == chunk_indices[probe_id] - sample_offset);
assert(sample_offset + list_length <= max_samples);
if constexpr (!kManageLocalTopK) {
// max_samples is zero/unused in the kManageLocalTopK mode
assert(sample_offset + list_length <= max_samples);
}

constexpr int kUnroll = raft::WarpSize / Veclen;
constexpr uint32_t kNumWarps = kThreadsPerBlock / raft::WarpSize;
Expand Down
5 changes: 4 additions & 1 deletion cpp/src/neighbors/ivf_flat/jit_lto_kernels/ivf_flat_interleaved_scan_kernel.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -833,7 +833,10 @@ __device__ __forceinline__ void interleaved_scan_kernel_impl(const uint32_t quer
uint32_t sample_offset = 0;
if (probe_id > 0) { sample_offset = chunk_indices[probe_id - 1]; }
assert(list_length == chunk_indices[probe_id] - sample_offset);
assert(sample_offset + list_length <= max_samples);
if constexpr (!kManageLocalTopK) {
// max_samples is zero/unused in the kManageLocalTopK mode
assert(sample_offset + list_length <= max_samples);
}

constexpr int kUnroll = raft::WarpSize / Veclen;
constexpr uint32_t kNumWarps = kThreadsPerBlock / raft::WarpSize;
Expand Down
Loading