Skip to content

Fix the IVF-PQ list codepacking loops to stride over the whole grid #2750

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 2 commits into from
Oct 8, 2026
+10 −6
Merged

Fix the IVF-PQ list codepacking loops to stride over the whole grid #2750

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
9291f7c
Fix the IVF-PQ list codepacking loops to stride over the whole grid
bdice Oct 5, 2026
154ba69
Merge branch 'main' into ivf-pq-codepacking-grid-stride
bdice Oct 7, 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
16 changes: 10 additions & 6 deletions cpp/src/neighbors/ivf_pq/ivf_pq_codepacking.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
* SPDX-License-Identifier: Apache-2.0
*/

Expand Down Expand Up @@ -228,7 +228,9 @@ __device__ void run_on_list(
uint32_t pq_dim,
Action action)
{
for (uint32_t ix = threadIdx.x + blockDim.x * blockIdx.x; ix < len; ix += blockDim.x) {
// Grid-stride loop: each vector is processed by exactly one thread.
for (uint32_t ix = threadIdx.x + blockDim.x * blockIdx.x; ix < len;
ix += blockDim.x * gridDim.x) {
const uint32_t src_ix = std::holds_alternative<uint32_t>(offset_or_indices)
? std::get<uint32_t>(offset_or_indices) + ix
: std::get<const uint32_t*>(offset_or_indices)[ix];
Expand All @@ -248,8 +250,9 @@ __device__ void write_list(
Action action)
{
using subwarp_align = raft::Pow2<SubWarpSize>;
uint32_t stride = subwarp_align::div(blockDim.x);
uint32_t ix = subwarp_align::div(threadIdx.x + blockDim.x * blockIdx.x);
// Grid-stride loop: each vector is processed by exactly one subwarp.
uint32_t stride = subwarp_align::div(blockDim.x) * gridDim.x;
uint32_t ix = subwarp_align::div(threadIdx.x + blockDim.x * blockIdx.x);
for (; ix < len; ix += stride) {
const uint32_t dst_ix = std::holds_alternative<uint32_t>(offset_or_indices)
? std::get<uint32_t>(offset_or_indices) + ix
Expand All @@ -268,8 +271,9 @@ __device__ void write_list_flat(uint8_t* out_flat_codes,
Action action)
{
using subwarp_align = raft::Pow2<SubWarpSize>;
uint32_t stride = subwarp_align::div(blockDim.x);
uint32_t ix = subwarp_align::div(threadIdx.x + blockDim.x * blockIdx.x);
// Grid-stride loop: each vector is processed by exactly one subwarp.
uint32_t stride = subwarp_align::div(blockDim.x) * gridDim.x;
uint32_t ix = subwarp_align::div(threadIdx.x + blockDim.x * blockIdx.x);
for (; ix < len; ix += stride) {
const uint32_t dst_ix = std::holds_alternative<uint32_t>(offset_or_indices)
? std::get<uint32_t>(offset_or_indices) + ix
Expand Down
Loading