[BUG]: Compute sanitizer memcheck failed with `UTILITY_TEST`
Open
Nobody has claimed this yet.
P0: Must have
type: bug
- Dominant language
- Cuda
- Stars
- 667
- Forks
- 120
- Avg merge
- 7d 5h
- Merged PRs (30d)
- 4
Description
Is this a duplicate?
- I confirmed there appear to be no duplicate issues for this bug (https://github.com/NVIDIA/cuCollections/issues)
Type of Bug
Silent Failure
Describe the bug
========= COMPUTE-SANITIZER
Randomness seeded to: 840931458
========= Invalid __global__ write of size 1 bytes
========= at void thrust::cuda_cub::assign_value<thrust::cuda_cub::tag, thrust::device_ptr<bool>, const bool *>(thrust::cuda_cub::execution_policy<T1> &, T2, T3)::war_nvbugs_881631::device_path(thrust::cuda_cub::execution_policy<thrust::cuda_cub::tag> &, thrust::device_ptr<bool>, const bool *)+0xd0 in /home/yunsongw/Work/cuCollections/build/_deps/cccl-src/thrust/thrust/system/cuda/detail/assign_value.h:58
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x7b3c42e0000a is out of bounds
========= and is 1 bytes after the nearest allocation at 0x7b3c42e00000 of size 10 bytes
========= Device Frame:void thrust::cuda_cub::assign_value<thrust::cuda_cub::tag, thrust::device_ptr<bool>, const bool *>(thrust::cuda_cub::execution_policy<T1> &, T2, T3)+0xd0 in /home/yunsongw/Work/cuCollections/build/_deps/cccl-src/thrust/thrust/system/cuda/detail/assign_value.h:62
========= Device Frame:void thrust::cuda_cub::assign_value<thrust::cuda_cub::tag, thrust::system::cpp::detail::tag, thrust::device_ptr<bool>, const bool *>(thrust::cuda_cub::cross_system<T1, T2> &, T3, T4)::war_nvbugs_881631::device_path(thrust::cuda_cub::cross_system<thrust::cuda_cub::tag, thrust::system::cpp::detail::tag> &, thrust::device_ptr<bool>, const bool *)+0xd0 in /home/yunsongw/Work/cuCollections/build/_deps/cccl-src/thrust/thrust/system/cuda/detail/assign_value.h:86
========= Device Frame:void thrust::cuda_cub::assign_value<thrust::cuda_cub::tag, thrust::system::cpp::detail::tag, thrust::device_ptr<bool>, const bool *>(thrust::cuda_cub::cross_system<T1, T2> &, T3, T4)+0xd0 in /home/yunsongw/Work/cuCollections/build/_deps/cccl-src/thrust/thrust/system/cuda/detail/assign_value.h:90
========= Device Frame:void thrust::reference<bool, thrust::device_ptr<bool>, thrust::device_reference<bool>>::strip_const_assign_value<thrust::cuda_cub::cross_system<thrust::cuda_cub::tag, thrust::system::cpp::detail::tag>, const bool *>(const T1 &, T2)+0xd0 in /home/yunsongw/Work/cuCollections/build/_deps/cccl-src/thrust/thrust/detail/reference.h:366
========= Device Frame:void thrust::reference<bool, thrust::device_ptr<bool>, thrust::device_reference<bool>>::assign_from<thrust::cuda_cub::tag, thrust::system::cpp::detail::tag, const bool *>(T1 *, T2 *, T3)+0xd0 in /home/yunsongw/Work/cuCollections/build/_deps/cccl-src/thrust/thrust/detail/reference.h:346
========= Device Frame:void thrust::reference<bool, thrust::device_ptr<bool>, thrust::device_reference<bool>>::assign_from<const bool *>(T1)+0xd0 in /home/yunsongw/Work/cuCollections/build/_deps/cccl-src/thrust/thrust/detail/reference.h:357
========= Device Frame:thrust::reference<bool, thrust::device_ptr<bool>, thrust::device_reference<bool>>::operator =(const bool &)+0xd0 in /home/yunsongw/Work/cuCollections/build/_deps/cccl-src/thrust/thrust/detail/reference.h:155
========= Device Frame:thrust::device_reference<bool>::operator =(const bool &)+0xd0 in /home/yunsongw/Work/cuCollections/build/_deps/cccl-src/thrust/thrust/device_reference.h:297
========= Device Frame:void check_hash_result_kernel_64<thrust::detail::normal_iterator<thrust::device_ptr<bool>>>(T1)+0xd0 in /home/yunsongw/Work/cuCollections/tests/utility/hash_test.cu:113
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: [0x334660]
========= in /lib/x86_64-linux-gnu/libcuda.so.1
========= Host Frame: [0x156e4]
========= in /usr/local/cuda-12.3/lib64/libcudart.so.12
========= Host Frame:cudaLaunchKernel [0x74631]
========= in /usr/local/cuda-12.3/lib64/libcudart.so.12
========= Host Frame:__device_stub__Z27check_hash_result_kernel_64IN6thrust6detail15normal_iteratorINS0_10device_ptrIbEEEEEvT_(thrust::detail::normal_iterator<thrust::device_ptr<bool> >&) in /tmp/tmpxft_00014bbd_00000000-6_hash_test.compute_90.cudafe1.stub.c:32 [0x1fd64]
========= in /home/yunsongw/Work/cuCollections/build/tests/./UTILITY_TEST
========= Host Frame:CATCH2_INTERNAL_TEST_4() in /home/yunsongw/Work/cuCollections/tests/utility/hash_test.cu:154 [0x250be]
========= in /home/yunsongw/Work/cuCollections/build/tests/./UTILITY_TEST
========= Host Frame:Catch::RunContext::invokeActiveTestCase() in src/catch2/internal/catch_run_context.cpp:536 [0xa164e]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::RunContext::runCurrentTest(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&) in src/catch2/internal/catch_run_context.cpp:498 [0xa20ed]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::RunContext::runTest(Catch::TestCaseHandle const&) in src/catch2/internal/catch_run_context.cpp:236 [0xa2ee2]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::Session::runInternal() in src/catch2/catch_session.cpp:332 [0x7ac0c]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::Session::run() in src/catch2/catch_session.cpp:263 [0x7af20]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:main in src/catch2/internal/catch_main.cpp:36 [0x11e0]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2Main.so.3.3.0
========= Host Frame:__libc_start_call_main in ../sysdeps/nptl/libc_start_call_main.h:58 [0x29d90]
========= in /lib/x86_64-linux-gnu/libc.so.6
========= Host Frame:__libc_start_main in ../csu/libc-start.c:379 [0x29e40]
========= in /lib/x86_64-linux-gnu/libc.so.6
========= Host Frame:_start [0xa065]
========= in /home/yunsongw/Work/cuCollections/build/tests/./UTILITY_TEST
=========
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaStreamSynchronize.
========= Saved host backtrace up to driver entry point at error
========= Host Frame: [0x480996]
========= in /lib/x86_64-linux-gnu/libcuda.so.1
========= Host Frame:cudaStreamSynchronize [0x743fb]
========= in /usr/local/cuda-12.3/lib64/libcudart.so.12
========= Host Frame:int cuco::test::count_if<thrust::detail::normal_iterator<thrust::device_ptr<bool> >, thrust::identity<bool> >(thrust::detail::normal_iterator<thrust::device_ptr<bool> >, thrust::detail::normal_iterator<thrust::device_ptr<bool> >, thrust::identity<bool>, CUstream_st*) [clone .isra.0] in /home/yunsongw/Work/cuCollections/tests/test_utils.hpp:54 [0x205f0]
========= in /home/yunsongw/Work/cuCollections/build/tests/./UTILITY_TEST
========= Host Frame:CATCH2_INTERNAL_TEST_4() in /home/yunsongw/Work/cuCollections/tests/test_utils.hpp:69 [0x2504d]
========= in /home/yunsongw/Work/cuCollections/build/tests/./UTILITY_TEST
========= Host Frame:Catch::RunContext::invokeActiveTestCase() in src/catch2/internal/catch_run_context.cpp:536 [0xa164e]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::RunContext::runCurrentTest(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&) in src/catch2/internal/catch_run_context.cpp:498 [0xa20ed]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::RunContext::runTest(Catch::TestCaseHandle const&) in src/catch2/internal/catch_run_context.cpp:236 [0xa2ee2]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::Session::runInternal() in src/catch2/catch_session.cpp:332 [0x7ac0c]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::Session::run() in src/catch2/catch_session.cpp:263 [0x7af20]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:main in src/catch2/internal/catch_main.cpp:36 [0x11e0]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2Main.so.3.3.0
========= Host Frame:__libc_start_call_main in ../sysdeps/nptl/libc_start_call_main.h:58 [0x29d90]
========= in /lib/x86_64-linux-gnu/libc.so.6
========= Host Frame:__libc_start_main in ../csu/libc-start.c:379 [0x29e40]
========= in /lib/x86_64-linux-gnu/libc.so.6
========= Host Frame:_start [0xa065]
========= in /home/yunsongw/Work/cuCollections/build/tests/./UTILITY_TEST
=========
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaGetLastError.
========= Saved host backtrace up to driver entry point at error
========= Host Frame: [0x480996]
========= in /lib/x86_64-linux-gnu/libcuda.so.1
========= Host Frame:cudaGetLastError [0x4cdc7]
========= in /usr/local/cuda-12.3/lib64/libcudart.so.12
========= Host Frame:int cuco::test::count_if<thrust::detail::normal_iterator<thrust::device_ptr<bool> >, thrust::identity<bool> >(thrust::detail::normal_iterator<thrust::device_ptr<bool> >, thrust::detail::normal_iterator<thrust::device_ptr<bool> >, thrust::identity<bool>, CUstream_st*) [clone .isra.0] in /home/yunsongw/Work/cuCollections/tests/test_utils.hpp:54 [0x20a9c]
========= in /home/yunsongw/Work/cuCollections/build/tests/./UTILITY_TEST
========= Host Frame:CATCH2_INTERNAL_TEST_4() in /home/yunsongw/Work/cuCollections/tests/test_utils.hpp:69 [0x2504d]
========= in /home/yunsongw/Work/cuCollections/build/tests/./UTILITY_TEST
========= Host Frame:Catch::RunContext::invokeActiveTestCase() in src/catch2/internal/catch_run_context.cpp:536 [0xa164e]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::RunContext::runCurrentTest(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&) in src/catch2/internal/catch_run_context.cpp:498 [0xa20ed]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::RunContext::runTest(Catch::TestCaseHandle const&) in src/catch2/internal/catch_run_context.cpp:236 [0xa2ee2]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::Session::runInternal() in src/catch2/catch_session.cpp:332 [0x7ac0c]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:Catch::Session::run() in src/catch2/catch_session.cpp:263 [0x7af20]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2.so.3.3.0
========= Host Frame:main in src/catch2/internal/catch_main.cpp:36 [0x11e0]
========= in /home/yunsongw/Work/cuCollections/build/_deps/catch2-build/src/libCatch2Main.so.3.3.0
========= Host Frame:__libc_start_call_main in ../sysdeps/nptl/libc_start_call_main.h:58 [0x29d90]
========= in /lib/x86_64-linux-gnu/libc.so.6
========= Host Frame:__libc_start_main in ../csu/libc-start.c:379 [0x29e40]
========= in /lib/x86_64-linux-gnu/libc.so.6
========= Host Frame:_start [0xa065]
========= in /home/yunsongw/Work/cuCollections/build/tests/./UTILITY_TEST
How to Reproduce
compute-sanitizer --tool memcheck ./UTILITY_TEST
Expected behavior
memcheck should pass without failures
Operating System
Ubuntu 22.04
nvidia-smi output
+---------------------------------------------------------------------------------------+
| NVIDIA-SMI 545.29.06 Driver Version: 545.29.06 CUDA Version: 12.3 |
|-----------------------------------------+----------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+======================+======================|
| 0 Quadro RTX 8000 Off | 00000000:17:00.0 Off | Off |
| 34% 35C P8 10W / 260W | 6MiB / 49152MiB | 0% Default |
| | | N/A |
+-----------------------------------------+----------------------+----------------------+
| 1 Quadro P620 Off | 00000000:B3:00.0 On | N/A |
| 36% 50C P0 N/A / N/A | 741MiB / 2048MiB | 1% Default |
| | | N/A |
+-----------------------------------------+----------------------+----------------------+
+---------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=======================================================================================|
| 0 N/A N/A 5071 G /usr/lib/xorg/Xorg 4MiB |
| 1 N/A N/A 5071 G /usr/lib/xorg/Xorg 288MiB |
| 1 N/A N/A 5280 G /usr/bin/gnome-shell 205MiB |
| 1 N/A N/A 8468 G ...seed-version=20240711-180158.427000 96MiB |
| 1 N/A N/A 29374 G ...yOnDemand --variations-seed-version 146MiB |
| 1 N/A N/A 63537 G /opt/cisco/anyconnect/bin/acwebhelper 0MiB |
+---------------------------------------------------------------------------------------+
NVCC version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Wed_Nov_22_10:17:15_PST_2023
Cuda compilation tools, release 12.3, V12.3.107
Build cuda_12.3.r12.3/compiler.33567101_0
Contributor guide
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start with tests/utility/hash_test.cu at lines 113 and 154, then inspect tests/test_utils.hpp at lines 54 and 69. Reproduce UTILITY_TEST under compute-sanitizer and trace the out-of-bounds write involving the boolean result buffer. Done means the test no longer reports the invalid write or launch failure.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- testing-qa
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 35/100