[BUG] cagra::search on vpq_f16_index triggers OOB for every search_algo (26.10 nightly)

Describe the bug

cuvs::neighbors::cagra::search on a cagra::vpq_f16_index<half, uint32_t> that was constructed by hand from (1) a graph produced by cagra::build(..., host_padded_dataset_view, attach_dataset_on_build=false) and (2) a device_vpq_dataset produced by preprocessing::quantize::pq::make_vpq_dataset(..., host_matrix_view) — the workflow the pq.hpp docstring documents as the CAGRA-Q pattern — triggers an out-of-bounds device read in the search kernel. Reproduces on all three search_algo variants (MULTI_CTA, SINGLE_CTA, MULTI_KERNEL).

Steps/Code to reproduce bug

Attached: repro_cuvs_vpq_search_oob.cu — single file, ~180 lines, only depends on libcuvs / libraft / librmm. 10K rows × 128 dim, pq_dim=32, pq_bits=8, graph_degree=32.

# Build (any recent nvcc; only public cuvs headers used)
nvcc -std=c++20 -O2 --extended-lambda --expt-relaxed-constexpr \
     -I$CONDA_PREFIX/include -I$CONDA_PREFIX/include/rapids \
     -I$CONDA_PREFIX/include/raft -I$CONDA_PREFIX/include/cuvs \
     -L$CONDA_PREFIX/lib -lcuvs -lcuvs_c -lrmm -lstdc++ \
     -DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE \
     -DRAFT_SYSTEM_LITTLE_ENDIAN=1 \
     -gencode arch=compute_89,code=sm_89 \
     repro_cuvs_vpq_search_oob.cu -o repro

# Run (immediate crash on MULTI_CTA search)
LD_LIBRARY_PATH=$CONDA_PREFIX/lib ./repro

Bare run output:

graph built: size=10000 dim=0 graph_degree=32
vpq trained
vpq_idx: size=10000 dim=128 graph_degree=32
--- searching with algo=MULTI_CTA ---
terminate called after throwing an instance of 'raft::cuda_error'
  what():  CUDA error encountered at:
    file=.../raft/core/interruptible.hpp line=294:
    call='query_result',
    Reason=cudaErrorIllegalAddress:an illegal memory access was encountered

Under compute-sanitizer --tool memcheck the failing read looks like:

========= Invalid __global__ read of size 4 bytes
=========     at search_multi_cta+0xd1f0
=========     by thread (0,0,0) in block (1,29,0)
=========     Access to 0x795d9e2d2d00 is out of bounds
=========     and is 240466944 bytes after the nearest allocation
=========     at 0x795d8fd7f100 of size 1 bytes
=========     Host Frame: cuvs::neighbors::cagra::detail::multi_cta_search::select_and_run<
=========                   __half, unsigned int, float, unsigned int,
=========                   cuvs::neighbors::filtering::none_sample_filter>(...)
=========     Host Frame: cuvs::neighbors::cagra::search<
=========                   __half, unsigned int,
=========                   dataset_view<vpq_dataset_container, __half, long, ...>,
=========                   unsigned int>(...)

The nearest allocation ... of size 1 bytes looks like a placeholder descriptor being passed as a real buffer.

Setting search_params::algo = SINGLE_CTA or MULTI_KERNEL produces the same class of cudaErrorIllegalAddress — different search kernel, same shape.

Expected behavior

cagra::search on a vpq_f16_index returns the top-K neighbours without a memory violation, matching the CAGRA-Q pattern documented at cuvs/preprocessing/quantize/pq.hpp:280:

// `idx` is a `cagra::index<float, uint32_t>` with graph built on dense rows.
// `padded` is a `device_padded_dataset_view<float, int64_t>` view of those same rows.
cuvs::neighbors::vpq_params vpq_params{};
auto vpq = cuvs::preprocessing::quantize::pq::make_vpq_dataset(res, vpq_params, padded.view());
idx.update_device_dataset_same_layout(res, vpq.as_dataset_view());

The reproducer takes the equivalent "construct a new vpq_f16_index from the graph + vpq view" shape because it's the natural path when attach_dataset_on_build = false was used at build time.

Environment details:

• Environment location: Bare-metal, Linux 6.17.0-1013-aws (Ubuntu 24.04), NVIDIA L40S (compute 8.9), driver 590.48.01, 491 GB host RAM.
• Method of cuVS install: conda (rapidsai-nightly).
• Relevant packages:

  libcuvs         26.10.00a206  cuda13_260819055025_2140532c  rapidsai-nightly
  libcuvs-headers 26.10.00a206  cuda13_260819055025_2140532c  rapidsai-nightly
  libraft         26.10.00a26   cuda13_260819053610_9aa17e57  rapidsai-nightly
  cuda-version    13.3          hcbadf70_3                    conda-forge
  

• Source tree HEAD matches the shipped .so: 2140532c ("Use librtcx for JIT-LTO", 2026-08-18).

Additional context

• Reproduces on every search_algo value (MULTI_CTA, SINGLE_CTA, MULTI_KERNEL), so it does not appear to be one specialised kernel — likely a mistake in the dataset-descriptor plumbing that all three search implementations consume.
• Traced from the pq.hpp:280 CAGRA-Q docstring example — the intended workflow.
• Superficially similar to issue #1802 but different kernel (multi_cta_search vs single_cta_search) and different dataset descriptor (vpq_dataset_container vs plain float). Fixing one may or may not cover this; worth checking whether the underlying dataset_descriptor_base_t construction path is shared.
• Build path itself is fine — cagra::build returns a host_padded_index with sane size() and graph(), and make_vpq_dataset returns a device_vpq_dataset with correct n_rows/dim/stride. The hand-stitched vpq_f16_index also reports correct size()/dim()/graph_degree(). Only the search kernel touches invalid memory.

Start by building and running the attached repro under compute-sanitizer, then inspect cagra::search and the dataset-descriptor construction used by all three search algorithms. Compare that path with the CAGRA-Q workflow documented at cuvs/preprocessing/quantize/pq.hpp:280 and issue #1802. Done means the repro returns top-K neighbours without an illegal device read for MULTI_CTA, SINGLE_CTA, or MULTI_KERNEL.