[BUG]: nullptr dereference UB triggered in for_each.h by thrust/testing/vector_allocators.cu test

[Artem-B]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Something else

Component

Thrust

Describe the bug

UB null sanitizer found another code path that leads to dereferencing nullptr and passing a nullptr reference around.

@miscco This may be another instance of #2964

cccl/thrust/thrust/system/detail/sequential/for_each.h

Lines 60 to 70 in 9019a6a

	_CCCL_HOST_DEVICE InputIterator
	for_each_n(sequential::execution_policy<DerivedPolicy>&, InputIterator first, Size n, UnaryFunction f)
	{
	// wrap f
	thrust::detail::wrapped_function<UnaryFunction, void> wrapped_f(f);
	for (Size i = 0; i != n; i++)
	{
	// we can dereference an OutputIterator if f does not
	// try to use the reference for anything besides assignment
	wrapped_f(*first);

It all starts here:

cccl/thrust/testing/vector_allocators.cu

Lines 187 to 190 in 9019a6a

	Vector v1(10, alloc1);
	Vector v2(15, alloc2);
	v2 = v1;

Full stack trace from cccl v2.6.1:

...third_party/gpus/cccl/v2_6_0/thrust/thrust/system/detail/sequential/for_each.h:70:15: runtime error: reference binding to null pointer of type 'int'
SUMMARY: UndefinedBehaviorSanitizer: undefined-behavior third_party/gpus/cccl/v2_6_0/thrust/thrust/system/detail/sequential/for_each.h:70:15

Thread 1 "vector_allocato" received signal SIGABRT, Aborted.
0x00007ffff7cd6981 in raise () from /usr/grte/v5/lib64/libc.so.6
(gdb) bt
#2  0x00005555592225af in __sanitizer::Abort () at third_party/llvm/llvm-project/compiler-rt/lib/sanitizer_common/sanitizer_posix_libcdep.cpp:163
#3  0x0000555559231dc1 in __sanitizer::Die () at third_party/llvm/llvm-project/compiler-rt/lib/sanitizer_common/sanitizer_termination.cpp:58
#4  0x000055555923932f in __ubsan_handle_type_mismatch_v1_abort () at third_party/llvm/llvm-project/compiler-rt/lib/ubsan/ubsan_handlers.cpp:148
#5  0x000055555925e9b0 in thrust::THRUST_200601_600_700_800_900_NS::system::detail::sequential::for_each_n<thrust::THRUST_200601_600_700_800_900_NS::system::cpp::detail::tag, int*, long, thrust::THRUST_200601_600_700_800_900_NS::detail::allocator_traits_detail::destroy_via_allocator<stateful_allocator<std::__u::allocator<int>, true> > > (first=0x0, n=15, f=...) at third_party/gpus/cccl/v2_6_0/thrust/thrust/system/detail/sequential/for_each.h:70
#6  0x000055555925e882 in thrust::THRUST_200601_600_700_800_900_NS::for_each_n<thrust::THRUST_200601_600_700_800_900_NS::system::cpp::detail::tag, int*, long, thrust::THRUST_200601_600_700_800_900_NS::detail::allocator_traits_detail::destroy_via_allocator<stateful_allocator<std::__u::allocator<int>, true> > > (exec=..., first=0x0, n=15, f=...) at third_party/gpus/cccl/v2_6_0/thrust/thrust/detail/for_each.inl:66
#7  0x000055555925e7ae in thrust::THRUST_200601_600_700_800_900_NS::detail::allocator_traits_detail::destroy_range<stateful_allocator<std::__u::allocator<int>, true>, int*, long> (a=..., p=0x0, n=15)
    at third_party/gpus/cccl/v2_6_0/thrust/thrust/detail/allocator/destroy_range.inl:98
#8  0x000055555925e5ca in thrust::THRUST_200601_600_700_800_900_NS::detail::destroy_range<stateful_allocator<std::__u::allocator<int>, true>, int*, long> (a=..., p=0x0, n=15)
    at third_party/gpus/cccl/v2_6_0/thrust/thrust/detail/allocator/destroy_range.inl:133
#9  0x000055555925dac9 in thrust::THRUST_200601_600_700_800_900_NS::detail::contiguous_storage<int, stateful_allocator<std::__u::allocator<int>, true> >::destroy (this=0x7fffffffc3f0, first=..., last=...)
    at third_party/gpus/cccl/v2_6_0/thrust/thrust/detail/contiguous_storage.inl:259
#10 0x00005555592a95a3 in thrust::THRUST_200601_600_700_800_900_NS::detail::vector_base<int, stateful_allocator<std::__u::allocator<int>, true> >::range_assign<thrust::THRUST_200601_600_700_800_900_NS::detail::normal_iterator<int const*> > (
    this=0x7fffffffc3f0, first=..., last=...) at third_party/gpus/cccl/v2_6_0/thrust/thrust/detail/vector_base.inl:999
#11 0x00005555592a93da in thrust::THRUST_200601_600_700_800_900_NS::detail::vector_base<int, stateful_allocator<std::__u::allocator<int>, true> >::range_assign<thrust::THRUST_200601_600_700_800_900_NS::detail::normal_iterator<int const*> > (
    this=0x7fffffffc3f0, first=..., last=...) at third_party/gpus/cccl/v2_6_0/thrust/thrust/detail/vector_base.inl:957
#12 0x00005555592a935a in thrust::THRUST_200601_600_700_800_900_NS::detail::vector_base<int, stateful_allocator<std::__u::allocator<int>, true> >::assign_dispatch<thrust::THRUST_200601_600_700_800_900_NS::detail::normal_iterator<int const*> > (
    this=0x7fffffffc3f0, first=..., last=...) at third_party/gpus/cccl/v2_6_0/thrust/thrust/detail/vector_base.inl:652
#13 0x00005555592a8e8a in thrust::THRUST_200601_600_700_800_900_NS::detail::vector_base<int, stateful_allocator<std::__u::allocator<int>, true> >::assign<thrust::THRUST_200601_600_700_800_900_NS::detail::normal_iterator<int const*> > (
    this=0x7fffffffc3f0, first=..., last=...) at third_party/gpus/cccl/v2_6_0/thrust/thrust/detail/vector_base.inl:607
#14 0x00005555592a8b9f in thrust::THRUST_200601_600_700_800_900_NS::detail::vector_base<int, stateful_allocator<std::__u::allocator<int>, true> >::operator= (this=0x7fffffffc3f0, v=...)
    at third_party/gpus/cccl/v2_6_0/thrust/thrust/detail/vector_base.inl:128
#15 0x00005555592a88dc in thrust::THRUST_200601_600_700_800_900_NS::host_vector<int, stateful_allocator<std::__u::allocator<int>, true> >::operator= (this=0x7fffffffc3f0, v=...) at third_party/gpus/cccl/v2_6_0/thrust/thrust/host_vector.h:163
#16 0x00005555592474cc in TestVectorAllocatorPropagateOnCopyAssignment<thrust::THRUST_200601_600_700_800_900_NS::host_vector<int, stateful_allocator<std::__u::allocator<int>, true> > > ()
    at third_party/gpus/cccl/v2_6_0/thrust/testing/vector_allocators.cu.cc:190
#17 0x0000555559240599 in TestVectorAllocatorPropagateOnCopyAssignmentHost () at third_party/gpus/cccl/v2_6_0/thrust/testing/vector_allocators.cu.cc:199
#18 0x000055555924d7e6 in TestVectorAllocatorPropagateOnCopyAssignmentHostUnitTest::run (this=0x55555b11c640 <TestVectorAllocatorPropagateOnCopyAssignmentHostInstance>) at third_party/gpus/cccl/v2_6_0/thrust/testing/vector_allocators.cu.cc:201
#19 0x00005555592b5365 in UnitTestDriver::run_tests (this=0x55555b11cef0 <driver_instance(thrust::THRUST_200601_600_700_800_900_NS::cuda_cub::tag)::s_instance>, tests_to_run=..., kwargs=...)
    at third_party/gpus/cccl/v2_6_0/thrust/testing/unittest/testframework.cu.cc:316
#20 0x00005555592b674a in UnitTestDriver::run_tests (this=0x55555b11cef0 <driver_instance(thrust::THRUST_200601_600_700_800_900_NS::cuda_cub::tag)::s_instance>, args=..., kwargs=...)
    at third_party/gpus/cccl/v2_6_0/thrust/testing/unittest/testframework.cu.cc:428
#21 0x00005555592d4ab6 in CUDATestDriver::run_tests (this=0x55555b11cef0 <driver_instance(thrust::THRUST_200601_600_700_800_900_NS::cuda_cub::tag)::s_instance>, args=..., kwargs=...)
    at third_party/gpus/cccl/v2_6_0/thrust/testing/unittest/cuda/testframework.cu.cc:205
#22 0x00005555592b74c8 in main (argc=1, argv=0x7fffffffd7c8) at third_party/gpus/cccl/v2_6_0/thrust/testing/unittest/testframework.cu.cc:530

How to Reproduce

Build and run thrust/testing/unittest/testframework.cu.cc w/o optimizations w/ clang and -fsanitize=null

Expected behavior

No undefinied behavior. checks for

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response