NVIDIA / NVIDIA/cccl

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

Open
#3,081 1 comment 0 reactions 0 assignees View on GitHub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

### Is this a duplicate?

- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### 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 https://github.com/NVIDIA/cccl/issues/2964

https://github.com/NVIDIA/cccl/blob/9019a6a28f2b6947d76c3b7502764a0a240f5111/thrust/thrust/system/detail/sequential/for_each.h#L60-L70

It all starts here:
https://github.com/NVIDIA/cccl/blob/9019a6a28f2b6947d76c3b7502764a0a240f5111/thrust/testing/vector_allocators.cu#L187-L190

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, 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, 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, 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, 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, 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, true> >::range_assign > (
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, true> >::range_assign > (
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, true> >::assign_dispatch > (
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, true> >::assign > (
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, 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, true> >::operator= (this=0x7fffffffc3f0, v=...) at third_party/gpus/cccl/v2_6_0/thrust/thrust/host_vector.h:163
#16 0x00005555592474cc in TestVectorAllocatorPropagateOnCopyAssignment, 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 ) at third_party/gpus/cccl/v2_6_0/thrust/testing/vector_allocators.cu.cc:201
#19 0x00005555592b5365 in UnitTestDriver::run_tests (this=0x55555b11cef0 , 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 , 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 , 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_

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.