cuvsCagraSearch returned 0 (Reason=cudaErrorInvalidValue:invalid argument)

[narangvivek10]

Issue:

[ERROR] com.nvidia.cuvs.lucene.TestCuVSRandomizedVectorSearch.testVectorSearch -- Time elapsed: 0.157 s <<< ERROR!
java.lang.RuntimeException: 
cuvsCagraSearch returned 0[CUDA error encountered at: file=/home/narang/code/cuvs/cpp/src/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh line=2207: call='cudaPeekAtLastError()', Reason=cudaErrorInvalidValue:invalid argument
Obtained 8 stack frames
#1 in /home/narang/code/cuvs/cpp/build/c/libcuvs_c.so: raft::cuda_error::cuda_error(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) +0x6e [0x77096bf4f90e]
#2 in /home/narang/code/cuvs/cpp/build/libcuvs.so(+0x1897edf) [0x77088f297edf]
#3 in /home/narang/code/cuvs/cpp/build/libcuvs.so: cuvs::neighbors::cagra::detail::single_cta_search::search<float, unsigned int, float, cuvs::neighbors::filtering::none_sample_filter, unsigned int, unsigned int>::operator()(raft::resources const&, cuda::std::__4::mdspan<unsigned int const, cuda::std::__4::extents<long, 18446744073709551615ul, 18446744073709551615ul>, cuda::std::__4::layout_right, raft::host_device_accessor<cuda::std::__4::default_accessor<unsigned int const>, (raft::memory_type)2> >, std::optional<cuda::std::__4::mdspan<unsigned int const, cuda::std::__4::extents<long, 18446744073709551615ul>, cuda::std::__4::layout_right, raft::host_device_accessor<cuda::std::__4::default_accessor<unsigned int const>, (raft::memory_type)2> > >, unsigned int*, float*, float const*, unsigned int, unsigned int const*, unsigned int*, unsigned int, cuvs::neighbors::filtering::none_sample_filter) +0x16f [0x77088f2df0bf]
#4 in /home/narang/code/cuvs/cpp/build/libcuvs.so(+0x18cc812) [0x77088f2cc812]
#5 in /home/narang/code/cuvs/cpp/build/libcuvs.so: void cuvs::neighbors::cagra::search<float, unsigned int, unsigned int>(raft::resources const&, cuvs::neighbors::cagra::search_params const&, cuvs::neighbors::cagra::index<float, unsigned int> const&, cuda::std::__4::mdspan<float const, cuda::std::__4::extents<long, 18446744073709551615ul, 18446744073709551615ul>, cuda::std::__4::layout_right, raft::host_device_accessor<cuda::std::__4::default_accessor<float const>, (raft::memory_type)2> >, cuda::std::__4::mdspan<unsigned int, cuda::std::__4::extents<long, 18446744073709551615ul, 18446744073709551615ul>, cuda::std::__4::layout_right, raft::host_device_accessor<cuda::std::__4::default_accessor<unsigned int>, (raft::memory_type)2> >, cuda::std::__4::mdspan<float, cuda::std::__4::extents<long, 18446744073709551615ul, 18446744073709551615ul>, cuda::std::__4::layout_right, raft::host_device_accessor<cuda::std::__4::default_accessor<float>, (raft::memory_type)2> >, cuvs::neighbors::filtering::base_filter const&) +0x3e5 [0x77088f2dbe15]
#6 in /home/narang/code/cuvs/cpp/build/c/libcuvs_c.so(+0xab40b) [0x77096bf8040b]
#7 in /home/narang/code/cuvs/cpp/build/c/libcuvs_c.so: cuvsCagraSearch +0x51 [0x77096bf817b1]
#8 in [0x7709983cf2e2]
]
	at __randomizedtesting.SeedInfo.seed([55C8B265D329C693:C1C1A0838FF56B00]:0)
	at com.nvidia.cuvs.internal.common.Util.checkCuVSError(Util.java:87)
	at com.nvidia.cuvs.internal.CagraIndexImpl.search(CagraIndexImpl.java:309)
	at com.nvidia.cuvs.lucene.CuVS2510GPUVectorsReader.search(CuVS2510GPUVectorsReader.java:453)
.
.
.

To reproduce this issue do:

mvn clean test -Dtest=TestCuVSRandomizedVectorSearch#testVectorSearch -Dtests.seed=55C8B265D329C693

[mythrocks]

There's a couple of things at play here, because I'm also seeing an error in device memory allocation:

java.lang.RuntimeException: cuvsCagraSearch returned 0[std::bad_alloc: out_of_memory: RMM failure at:/home/mithunr/workspace/dev/cuvs/1/cpp/build/_deps/rmm-src/cpp/include/rmm/mr/limiting_resource_adaptor.hpp:148: Exceeded memory limit (failed to allocate 4.000000 B)]
        at __randomizedtesting.SeedInfo.seed([55C8B265D329C693:C1C1A0838FF56B00]:0)
        at com.nvidia.cuvs.internal.common.Util.checkCuVSError(Util.java:87)
        at com.nvidia.cuvs.internal.CagraIndexImpl.search(CagraIndexImpl.java:309)
        at com.nvidia.cuvs.lucene.CuVS2510GPUVectorsReader.search(CuVS2510GPUVectorsReader.java:447)
        at org.apache.lucene.codecs.perfield.PerFieldKnnVectorsFormat$FieldsReader.search(PerFieldKnnVectorsFormat.java:311)
        at org.apache.lucene.index.CodecReader.searchNearestVectors(CodecReader.java:272)
        at org.apache.lucene.search.KnnFloatVectorQuery.approximateSearch(KnnFloatVectorQuery.java:119)
        at org.apache.lucene.search.AbstractKnnVectorQuery.getLeafResults(AbstractKnnVectorQuery.java:131)
        at org.apache.lucene.search.AbstractKnnVectorQuery.searchLeaf(AbstractKnnVectorQuery.java:113)
        at org.apache.lucene.search.AbstractKnnVectorQuery.lambda$rewrite$0(AbstractKnnVectorQuery.java:96)
        at java.base/java.util.concurrent.FutureTask.run(FutureTask.java:317)
        at org.apache.lucene.search.TaskExecutor$Task.run(TaskExecutor.java:173)
        at org.apache.lucene.search.TaskExecutor.invokeAll(TaskExecutor.java:111)
        at org.apache.lucene.search.AbstractKnnVectorQuery.rewrite(AbstractKnnVectorQuery.java:98)
        at org.apache.lucene.search.KnnFloatVectorQuery.rewrite(KnnFloatVectorQuery.java:49)
        at org.apache.lucene.search.IndexSearcher.rewrite(IndexSearcher.java:874)
        at org.apache.lucene.tests.search.AssertingIndexSearcher.rewrite(AssertingIndexSearcher.java:67)
        at org.apache.lucene.search.IndexSearcher.rewrite(IndexSearcher.java:885)
        at org.apache.lucene.search.IndexSearcher.search(IndexSearcher.java:748)
        at org.apache.lucene.search.IndexSearcher.searchAfter(IndexSearcher.java:596)
        at org.apache.lucene.search.IndexSearcher.search(IndexSearcher.java:619)
        at com.nvidia.cuvs.lucene.TestCuVSRandomizedVectorSearch.testVectorSearch(TestCuVSRandomizedVectorSearch.java:130)
...

It looks like the CUDA kernel error (from the description) appears first, and then the next device memory allocation fails.

[mythrocks]

I should also document that compute-sanitizer reveals that a couple of other problems that might need addressing. I'm not sure if they're related yet:

Init-checks:

========= Uninitialized __global__ memory read of size 4 bytes
=========     at const T1 * cuvs::neighbors::cagra::detail::setup_workspace_standard<cuvs::neighbors::cagra::detail::standard_dataset_descriptor_t<(cuvs::distance::DistanceType)0, (unsigned int)16, (unsigned int)256, float, unsigned int, float>>(const T1 *, void *, const T1::DATA_T *, unsigned int)+0x110
=========     by thread (12,0,0) in block (0,0,0)
=========     Address 0x14be000030
=========         Device Frame: 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)+0x170
=========     Saved host backtrace up to driver entry point at kernel launch time
=========         Host Frame: cuLaunchKernel_ptsz [0x3526e4] in libcuda.so.1
=========         Host Frame:  [0x143a8] in libcudart.so.12.9.79
=========         Host Frame: cudaLaunchKernel_ptsz [0x5735d] in libcudart.so.12.9.79
...

Race-checks:

========= Warning: Race reported between Write access 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)+0x38d0
=========     and Read access 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)+0x50d0 [16 hazards]
=========

cc @benfred, @cjnolet, to keep them apprised.

[mythrocks]

For the record:

1. The RMM failure indicated above was largely resolved with Setup ThreadLocal based CuVSResources usage #87.
2. I strongly suspect that the kernel failure in the description is the result of a possible error in synchronization, in the Java layer.

When tracing through the CagraIndexImpl (in cuvs-java), one sees that there are two threads at play during the crash:

1. The main thread, whose query object carries the same CuVSResources as the CagraIndex does.
2. Another thread which seems to have its own thread-local CuVSResources. I'm wondering if this is the reason for the failure: that the index is being queried on separate threads, and synching on different cuda-streams.

Still investigating.

[mythrocks]

Pasting advice from @viclafargue here, with much gratitude:

I believe that the issue comes from running multiple CAGRA searches in parallel without using the dedicated persistent kernel feature.
Modern GPUs support much larger use of shared memory, but the user have to "opt-in" by setting the cudaFuncAttributeMaxDynamicSharedMemorySize attribute of a kernel. This is what is being done in the normal path (not persistent) of CAGRA search.
The issue is that different types of queries require varying amount of shared memory for the search kernel. And the global attribute of this same kernel is set to different values in different threads. When reaching the kernel launch an invalid amount of shared memory is requested (larger than attribute value).
I think there are multiple solutions :

1. making sure to never perform searches in parallel by syncing on the same stream
2. using persistent kernels
3. or maybe fixing this by setting the kernel's attribute once in a thread safe fashion (std::call_once) to set it to sharedMemPerBlockOptin.

Sharing the same cuvsResources across multiple threads would allow us to synchronize on the same cuda-stream. Per advice from @cjnolet, sharing cuvsResources isn't an option.

I tried introducing synchronization by sharing just the stream across threads. It didn't seem to resolve the issue on the first go. I'm still working on trying that.

Switching to a persistent kernel: I'll leave that as a last resort. It will need examination by @benfred, @cjnolet.

[cjnolet]

The problem with the persistent kernel is that it's limited by k, unfortunately.

If what @viclafargue is saying is true, the I think we need to rethink these kernels a bit.

Also tagging @tfeher and @achirkin (and @anaruse @enp1s0) for additonal thoughts.

[viclafargue]

making sure to never perform searches in parallel by syncing on the same stream

Actually, after thinking a bit more about it, I am not sure that syncing on the same CUDA stream would be sufficient. The issue is that the cudaFuncAttributeMaxDynamicSharedMemorySize attribute of a CUDA kernel is a global value tied to the CUDA context used by all the threads. Threads overwrite each others when setting its value.

Here are updated 3 solutions :

1. Ensuring that CAGRA searches are never done in parallel by never using multi-threading or by protecting the specific case of CAGRA search with a mutex.
2. Using the persistent kernel feature (see the parameters over here). This option keeps the search kernel alive for a few seconds while the work is being queued in pinned memory, it offers the best performance for series of small queries.
3. Fixing the cuVS code that sets cudaFuncAttributeMaxDynamicSharedMemorySize to make it compatible with multi-threading. But, there might be other hidden issues that would prevent multi-threading from working.

[mythrocks]

Thank you, @viclafargue. I should have a patch up for this tomorrow. Thank you again for the advice.

[mythrocks]

For the record, I've posted rapidsai/cuvs#1771 to address this failure.