[BUG]: Compute sanitizer memcheck failed with UTILITY_TEST

[PointKernel]

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