NVIDIA / NVIDIA/cuvs

[BUG] OOB access in multi-threaded CAGRA search

Open
#1,802 0 comments 0 reactions 1 assignee View on GitHub

@viclafargue is already working on this.

Since Feb 13, 2026.

bug
Dominant language
Cuda
Stars
854
Forks
236
Avg merge
3d 3h
Merged PRs (30d)
62

Description

========= Invalid __global__ read of size 8 bytes
=========     at void cuvs::neighbors::cagra::detail::single_cta_search::search_kernel<(unsigned int)64, (unsigned int)64, (unsigned int)1, cuvs::neighbors::cagra::detail::dataset_descriptor_base_t<float, unsigned int, float>, unsigned int, cuvs::neighbors::filtering::none_sample_filter>(unsigned long, T4::DISTANCE_T *, unsigned int, const T4*, const T4::DATA_T *, const T4::INDEX_T *, unsigned int, const T5 *, unsigned int, unsigned long, const T4::INDEX_T *, unsigned int, T4::INDEX_T *, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int *, unsigned int, unsigned int, unsigned int, T6)+0x40
=========     by thread (0,0,0) in block (0,15,0)
=========     Access to 0x6290000020 is potentially made before memory is allocated
=========     and is inside the nearest allocation at 0x6290000000 of size 64 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========         Host Frame: cuLaunchKernel_ptsz [0x3534c4] in libcuda.so.1
=========         Host Frame:  [0x14e7c] in libcudart.so.13
=========         Host Frame: cudaLaunchKernel_ptsz [0x572b7] in libcudart.so.13

Contributor guide

Open the contributing guide

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.