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

[bdice]

write_list, write_list_flat and run_on_list in ivf_pq_codepacking.cuh started at the global row index but advanced by one block's worth of rows:

uint32_t stride = subwarp_align::div(blockDim.x);                             // rows per block
uint32_t ix     = subwarp_align::div(threadIdx.x + blockDim.x * blockIdx.x);  // global start
for (; ix < len; ix += stride) { ... }

Their launchers already size the grid to cover every row (encode_list_data launches one block per 8 rows, or 16 for 4-bit codes; pack/unpack/reconstruct launch one block per 256 rows), so block b processed every row from its first one to the end of the list. Rows were encoded or packed up to n_rows / rows_per_block times, and the subwarps of block 0 walked the entire list, each row a serial loop over all pq_dim subspaces. For the ~128-row lists in NEIGHBORS_ANN_IVF_PQ_TEST that is up to 16 passes per row; at pq_dim = 3072 one encode_list_data_interleaved_kernel launch took ~127–146 ms.

Affected paths:

• encode_list_data, used by ivf_pq::helpers::codepacker::extend_list, for any list longer than one block (8–16 rows).
• pack_list_data, unpack_list_data, {un}pack_contiguous_list_data and reconstruct_list_data for lists longer than 256 rows. The C API (cuvsIvfPqIndexUnpackContiguousListData) and Python index.lists() go through these; unpacking a 100k-row list did ~2·10⁷ row-unpacks instead of 10⁵.

The main build / extend encoder (process_and_fill_codes_kernel) does not use these loops and is unaffected.

This PR makes the three loops grid-stride loops by multiplying the stride by gridDim.x. Launch configurations are unchanged.

• Codes are identical. Each row is still handled by the same per-row routine (encode_vectors, or the pack/unpack/reconstruct action) with the same lane layout, codebook scan order and reduction order, and is written to the same location. The only difference is that each row is now processed exactly once instead of up to 16 times.
• This also removes a possible lost update: with pq_bits < 8, the FLAT-layout encoder and unpack_contiguous do bitfield read-modify-writes on shared bytes, and two blocks processing the same row could collide.

The commit also updates the file's SPDX copyright header to the canonical form required by the copyright pre-commit hook.

Testing

All IVF-PQ test executables and the full ctest suite pass. The existing tests already check that packing round-trips byte-exactly and that reconstruct → extend reproduces the vectors.

Measurements

Single-process wall time on an RTX 6000 Ada (48 GB), otherwise idle (no ctest parallelism). The old and new libcuvs.so were run alternately, 2 or more repetitions each, with the order reversed between repetitions; mean (min–max). The two builds differ only by this change.

executable	before	after	change
NEIGHBORS_ANN_IVF_PQ_TEST	158.6 s (157.9–159.4)	138.5 s (138.2–138.7)	-12.7%

The other executables measured the same way (IVF-Flat, IVF-SQ, CAGRA, multi-GPU, tiered index, dynamic batching, all-neighbors) were unchanged within noise. All tests passed in every run.

[copy-pr-bot]

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.