Hacktoberfest 2026: le issue che i maintainer hanno segnato per ottobre, aperte e adatte ai principianti. Sfoglia le issue Hacktoberfest

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

Aperta
  • Qualcosa non va?
#2,489 1 commento 0 reazioni 0 assegnatari Vedi su GitHub

I maintainer di solito rispondono entro 1 giorno

Nessuno ha ancora preso questa issue.

Valutazione

Difficoltà
4/5
Tempo stimato
3-5 giorni
Idoneità per principianti
48/100
Tipo di issue
Bug
Chiarezza
Abbastanza chiara
Stato di attività
Tranquilla
Stack tecnologico
cpp
Ambito
search

Direzione di ricerca

Inizia compilando ed eseguendo il repro allegato con compute-sanitizer, quindi ispeziona cagra::search e la costruzione del descrittore del dataset utilizzata da tutti e tre gli algoritmi di ricerca. Confronta questo percorso con il workflow CAGRA-Q documentato in cuvs/preprocessing/quantize/pq.hpp:280 e in issue #1802. Il lavoro è completo quando il repro restituisce i vicini top-K senza una lettura illegale dal device per MULTI_CTA, SINGLE_CTA o MULTI_KERNEL.

Scritto dal modello di indicizzazione a partire dal testo della issue.

Descrizione

bug

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.
Lingua principale
Cuda
Stelle
863
Fork
242
Merge medio
2g 4h
PR unite (30g)
80

Preparare l'ambiente

Come iniziare

  1. Leggi tutta la issue e poi la guida ai contributi del progetto.
  2. Commenta sulla issue per dire che te ne occupi tu — evita che due persone facciano lo stesso lavoro.
  3. Fai un fork del repository e lavora su un branch.
  4. Apri una pull request che faccia riferimento al numero della issue.

Altre issue di NVIDIA/cuvs

Tutte le issue di NVIDIA/cuvs

Issue simili

Altre issue su Search

Ricevi le nuove issue nella tua casella

Un breve riepilogo di issue GitHub adatte ai principianti.