Bug in thrust::system::detail::sequential::copy_detail::copy for trivially copyable types

[Eskilade]

Errors reported by clangs UBSan on this example :

#include <thrust/host_vector.h>
int main()
{
 thrust::host_vector<int> v(0);
 v.resize(1);
 return 0;
}

Output :

[user@ez cuda]$ ./a.out 
/usr/local/cuda-9.2/include/thrust/iterator/iterator_adaptor.h:191:14: runtime error: reference binding to null pointer of type 'int'
    #0 0x42db87 in thrust::iterator_adaptor<thrust::detail::normal_iterator<int*>, int*, thrust::use_default, thrust::use_default, thrust::use_default, thrust::use_default, thrust::use_default>::dereference() const /usr/local/cuda-9.2/include/thrust/iterator/iterator_adaptor.h:191:7
    #1 0x42d9df in thrust::detail::normal_iterator<int*>::reference thrust::iterator_core_access::dereference<thrust::detail::normal_iterator<int*> >(thrust::detail::normal_iterator<int*> const&) /usr/local/cuda-9.2/include/thrust/iterator/iterator_facade.h:128:16
    #2 0x42d5ab in thrust::iterator_facade<thrust::detail::normal_iterator<int*>, int, thrust::system::cpp::detail::tag, thrust::random_access_traversal_tag, int&, long>::operator*() const /usr/local/cuda-9.2/include/thrust/iterator/iterator_facade.h:310:14
    #3 0x42ea24 in int* thrust::system::detail::sequential::copy_detail::copy<thrust::detail::normal_iterator<int*>, int*>(thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*, thrust::detail::integral_constant<bool, true>) /usr/local/cuda-9.2/include/thrust/system/detail/sequential/copy.inl:60:59
    #4 0x42e9d4 in int* thrust::system::detail::sequential::copy<thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::system::detail::sequential::execution_policy<thrust::system::cpp::detail::tag>&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/system/detail/sequential/copy.inl:120:10
    #5 0x42e8ca in int* thrust::copy<thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::detail::execution_policy_base<thrust::system::cpp::detail::tag> const&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/copy.inl:37:10
    #6 0x42e7f1 in int* thrust::detail::two_system_copy<thrust::system::cpp::detail::tag, thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::execution_policy<thrust::system::cpp::detail::tag> const&, thrust::execution_policy<thrust::system::cpp::detail::tag> const&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/copy.inl:72:10
    #7 0x42e664 in _ZN6thrust6detail23allocator_traits_detail20copy_construct_rangeINS_6system3cpp6detail3tagESaIiENS0_15normal_iteratorIPiEES9_EENS0_10disable_ifIXsr34needs_copy_construct_via_allocatorIT0_NS0_15pointer_elementIT2_E4typeEEE5valueESE_E4typeERNS_16execution_policyIT_EERSC_T1_SO_SE_ /usr/local/cuda-9.2/include/thrust/detail/allocator/copy_construct_range.inl:218:10
    #8 0x42e5a0 in int* thrust::detail::copy_construct_range<thrust::system::cpp::detail::tag, std::allocator<int>, thrust::detail::normal_iterator<int*>, int*>(thrust::execution_policy<thrust::system::cpp::detail::tag>&, std::allocator<int>&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/allocator/copy_construct_range.inl:291:10
    #9 0x42e202 in thrust::detail::normal_iterator<int*> thrust::detail::contiguous_storage<int, std::allocator<int> >::uninitialized_copy<thrust::detail::normal_iterator<int*> >(thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>) /usr/local/cuda-9.2/include/thrust/detail/contiguous_storage.inl:225:19
    #10 0x42c8ac in thrust::detail::vector_base<int, std::allocator<int> >::append(unsigned long) /usr/local/cuda-9.2/include/thrust/detail/vector_base.inl:813:28
    #11 0x4295eb in thrust::detail::vector_base<int, std::allocator<int> >::resize(unsigned long) /usr/local/cuda-9.2/include/thrust/detail/vector_base.inl:261:5
    #12 0x428dfe in main /home/user/tests/cuda/main.cpp:6:4
    #13 0x7ff186e2afe9 in __libc_start_main (/lib64/libc.so.6+0x20fe9)
    #14 0x402d69 in _start (/home/user/tests/cuda/a.out+0x402d69)

/usr/local/cuda-9.2/include/thrust/iterator/iterator_facade.h:128:16: runtime error: reference binding to null pointer of type 'int'
    #0 0x42da19 in thrust::detail::normal_iterator<int*>::reference thrust::iterator_core_access::dereference<thrust::detail::normal_iterator<int*> >(thrust::detail::normal_iterator<int*> const&) /usr/local/cuda-9.2/include/thrust/iterator/iterator_facade.h:128:7
    #1 0x42d5ab in thrust::iterator_facade<thrust::detail::normal_iterator<int*>, int, thrust::system::cpp::detail::tag, thrust::random_access_traversal_tag, int&, long>::operator*() const /usr/local/cuda-9.2/include/thrust/iterator/iterator_facade.h:310:14
    #2 0x42ea24 in int* thrust::system::detail::sequential::copy_detail::copy<thrust::detail::normal_iterator<int*>, int*>(thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*, thrust::detail::integral_constant<bool, true>) /usr/local/cuda-9.2/include/thrust/system/detail/sequential/copy.inl:60:59
    #3 0x42e9d4 in int* thrust::system::detail::sequential::copy<thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::system::detail::sequential::execution_policy<thrust::system::cpp::detail::tag>&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/system/detail/sequential/copy.inl:120:10
    #4 0x42e8ca in int* thrust::copy<thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::detail::execution_policy_base<thrust::system::cpp::detail::tag> const&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/copy.inl:37:10
    #5 0x42e7f1 in int* thrust::detail::two_system_copy<thrust::system::cpp::detail::tag, thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::execution_policy<thrust::system::cpp::detail::tag> const&, thrust::execution_policy<thrust::system::cpp::detail::tag> const&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/copy.inl:72:10
    #6 0x42e664 in _ZN6thrust6detail23allocator_traits_detail20copy_construct_rangeINS_6system3cpp6detail3tagESaIiENS0_15normal_iteratorIPiEES9_EENS0_10disable_ifIXsr34needs_copy_construct_via_allocatorIT0_NS0_15pointer_elementIT2_E4typeEEE5valueESE_E4typeERNS_16execution_policyIT_EERSC_T1_SO_SE_ /usr/local/cuda-9.2/include/thrust/detail/allocator/copy_construct_range.inl:218:10
    #7 0x42e5a0 in int* thrust::detail::copy_construct_range<thrust::system::cpp::detail::tag, std::allocator<int>, thrust::detail::normal_iterator<int*>, int*>(thrust::execution_policy<thrust::system::cpp::detail::tag>&, std::allocator<int>&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/allocator/copy_construct_range.inl:291:10
    #8 0x42e202 in thrust::detail::normal_iterator<int*> thrust::detail::contiguous_storage<int, std::allocator<int> >::uninitialized_copy<thrust::detail::normal_iterator<int*> >(thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>) /usr/local/cuda-9.2/include/thrust/detail/contiguous_storage.inl:225:19
    #9 0x42c8ac in thrust::detail::vector_base<int, std::allocator<int> >::append(unsigned long) /usr/local/cuda-9.2/include/thrust/detail/vector_base.inl:813:28
    #10 0x4295eb in thrust::detail::vector_base<int, std::allocator<int> >::resize(unsigned long) /usr/local/cuda-9.2/include/thrust/detail/vector_base.inl:261:5
    #11 0x428dfe in main /home/user/tests/cuda/main.cpp:6:4
    #12 0x7ff186e2afe9 in __libc_start_main (/lib64/libc.so.6+0x20fe9)
    #13 0x402d69 in _start (/home/user/tests/cuda/a.out+0x402d69)

/usr/local/cuda-9.2/include/thrust/iterator/iterator_facade.h:310:14: runtime error: reference binding to null pointer of type 'int'
    #0 0x42d5e5 in thrust::iterator_facade<thrust::detail::normal_iterator<int*>, int, thrust::system::cpp::detail::tag, thrust::random_access_traversal_tag, int&, long>::operator*() const /usr/local/cuda-9.2/include/thrust/iterator/iterator_facade.h:310:7
    #1 0x42ea24 in int* thrust::system::detail::sequential::copy_detail::copy<thrust::detail::normal_iterator<int*>, int*>(thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*, thrust::detail::integral_constant<bool, true>) /usr/local/cuda-9.2/include/thrust/system/detail/sequential/copy.inl:60:59
    #2 0x42e9d4 in int* thrust::system::detail::sequential::copy<thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::system::detail::sequential::execution_policy<thrust::system::cpp::detail::tag>&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/system/detail/sequential/copy.inl:120:10
    #3 0x42e8ca in int* thrust::copy<thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::detail::execution_policy_base<thrust::system::cpp::detail::tag> const&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/copy.inl:37:10
    #4 0x42e7f1 in int* thrust::detail::two_system_copy<thrust::system::cpp::detail::tag, thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::execution_policy<thrust::system::cpp::detail::tag> const&, thrust::execution_policy<thrust::system::cpp::detail::tag> const&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/copy.inl:72:10
    #5 0x42e664 in _ZN6thrust6detail23allocator_traits_detail20copy_construct_rangeINS_6system3cpp6detail3tagESaIiENS0_15normal_iteratorIPiEES9_EENS0_10disable_ifIXsr34needs_copy_construct_via_allocatorIT0_NS0_15pointer_elementIT2_E4typeEEE5valueESE_E4typeERNS_16execution_policyIT_EERSC_T1_SO_SE_ /usr/local/cuda-9.2/include/thrust/detail/allocator/copy_construct_range.inl:218:10
    #6 0x42e5a0 in int* thrust::detail::copy_construct_range<thrust::system::cpp::detail::tag, std::allocator<int>, thrust::detail::normal_iterator<int*>, int*>(thrust::execution_policy<thrust::system::cpp::detail::tag>&, std::allocator<int>&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/allocator/copy_construct_range.inl:291:10
    #7 0x42e202 in thrust::detail::normal_iterator<int*> thrust::detail::contiguous_storage<int, std::allocator<int> >::uninitialized_copy<thrust::detail::normal_iterator<int*> >(thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>) /usr/local/cuda-9.2/include/thrust/detail/contiguous_storage.inl:225:19
    #8 0x42c8ac in thrust::detail::vector_base<int, std::allocator<int> >::append(unsigned long) /usr/local/cuda-9.2/include/thrust/detail/vector_base.inl:813:28
    #9 0x4295eb in thrust::detail::vector_base<int, std::allocator<int> >::resize(unsigned long) /usr/local/cuda-9.2/include/thrust/detail/vector_base.inl:261:5
    #10 0x428dfe in main /home/user/tests/cuda/main.cpp:6:4
    #11 0x7ff186e2afe9 in __libc_start_main (/lib64/libc.so.6+0x20fe9)
    #12 0x402d69 in _start (/home/user/tests/cuda/a.out+0x402d69)

/usr/local/cuda-9.2/include/thrust/system/detail/sequential/trivial_copy.h:44:24: runtime error: null pointer passed as argument 2, which is declared to never be null
/usr/include/string.h:47:14: note: nonnull attribute specified here
    #0 0x42eb78 in int* thrust::system::detail::sequential::trivial_copy_n<int>(int const*, long, int*) /usr/local/cuda-9.2/include/thrust/system/detail/sequential/trivial_copy.h:44:3
    #1 0x42ea51 in int* thrust::system::detail::sequential::copy_detail::copy<thrust::detail::normal_iterator<int*>, int*>(thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*, thrust::detail::integral_constant<bool, true>) /usr/local/cuda-9.2/include/thrust/system/detail/sequential/copy.inl:60:3
    #2 0x42e9d4 in int* thrust::system::detail::sequential::copy<thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::system::detail::sequential::execution_policy<thrust::system::cpp::detail::tag>&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/system/detail/sequential/copy.inl:120:10
    #3 0x42e8ca in int* thrust::copy<thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::detail::execution_policy_base<thrust::system::cpp::detail::tag> const&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/copy.inl:37:10
    #4 0x42e7f1 in int* thrust::detail::two_system_copy<thrust::system::cpp::detail::tag, thrust::system::cpp::detail::tag, thrust::detail::normal_iterator<int*>, int*>(thrust::execution_policy<thrust::system::cpp::detail::tag> const&, thrust::execution_policy<thrust::system::cpp::detail::tag> const&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/copy.inl:72:10
    #5 0x42e664 in _ZN6thrust6detail23allocator_traits_detail20copy_construct_rangeINS_6system3cpp6detail3tagESaIiENS0_15normal_iteratorIPiEES9_EENS0_10disable_ifIXsr34needs_copy_construct_via_allocatorIT0_NS0_15pointer_elementIT2_E4typeEEE5valueESE_E4typeERNS_16execution_policyIT_EERSC_T1_SO_SE_ /usr/local/cuda-9.2/include/thrust/detail/allocator/copy_construct_range.inl:218:10
    #6 0x42e5a0 in int* thrust::detail::copy_construct_range<thrust::system::cpp::detail::tag, std::allocator<int>, thrust::detail::normal_iterator<int*>, int*>(thrust::execution_policy<thrust::system::cpp::detail::tag>&, std::allocator<int>&, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, int*) /usr/local/cuda-9.2/include/thrust/detail/allocator/copy_construct_range.inl:291:10
    #7 0x42e202 in thrust::detail::normal_iterator<int*> thrust::detail::contiguous_storage<int, std::allocator<int> >::uninitialized_copy<thrust::detail::normal_iterator<int*> >(thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>, thrust::detail::normal_iterator<int*>) /usr/local/cuda-9.2/include/thrust/detail/contiguous_storage.inl:225:19
    #8 0x42c8ac in thrust::detail::vector_base<int, std::allocator<int> >::append(unsigned long) /usr/local/cuda-9.2/include/thrust/detail/vector_base.inl:813:28
    #9 0x4295eb in thrust::detail::vector_base<int, std::allocator<int> >::resize(unsigned long) /usr/local/cuda-9.2/include/thrust/detail/vector_base.inl:261:5
    #10 0x428dfe in main /home/user/tests/cuda/main.cpp:6:4
    #11 0x7ff186e2afe9 in __libc_start_main (/lib64/libc.so.6+0x20fe9)
    #12 0x402d69 in _start (/home/user/tests/cuda/a.out+0x402d69)

[dawagner]

I strongly believe that the fix should be done in the vector_base::append code and not somewhere deeper in the stack: once vector_base::append calls

m_storage.uninitialized_copy(begin(), end(), new_storage.begin());

the information whether the iterators are valid or not is lost: since neither the this pointer nor the new_storage object are passed, it is impossible to write something like x == this->end() or x.is_valid() or worse, (&*x) == NULL (see below) (where x is any of the 3 arguments passed to uninitialized_copy).

As a consequence, I think #938 was the correct fix and should be reopened.

Also, @griwes: you mentioned that there is a discussion for whether calling thrust::system::detail::sequential::trivial_copy_n(get(&*first), n, get(&*result)) may or may not result in UB; it seems obvious to me that it is if first can't be dereferenced. It seems that in the particular callstack pasted above, UBSan complains about the binding a reference to a null pointer and then about this null pointer being passed as argument to memmove which expects the source argument to be nonnull.

In a unit test that encounters this UB, we got the following crash:

terminate called after throwing an instance of 'thrust::system::system_error'

 what():  device free failed: an illegal memory access was encountered

At the time of writing, I have not thoroughly checked that this is related but I'm willing to bet it is.

[griwes]

There is one or two different UBs that are happening here, depending on which side in the debate I mentioned you take:

• &*first may dereference a null pointer (some dispute that this is UB, and UBSan is neither a final authority nor even an argument here), and
• memmove is called with a null pointer, which is illegal even when n is 0 (two points here: first, C is insane for defining this to be UB; and this is a different kind of UB than the first one, because it's "soft", library UB, not hard, core language UB, but that's besides the point).

Putting aside a problem that exists in all cases (that begin() and end() do not form valid pointers into the same array; this is a consequence of the fact that the vector elements are constructed separately, and not as an array, and is a known reason for why vector can't actually be implemented without UB), the call to uninitialized_copy is perfectly fine; the algorithm itself should do nothing when begin == end (this is the basis of how the STL algorithms have worked forever, as they operate on [begin, end), not [begin, end]).

The reson why the fix both can and must be applied down the stack is that the calls above are 100% correct, it's the backend that introduces the illegal behavior, and only in the trivial case; and also it is trivial to check, by guarding the call to trivial_copy_n against this case (either by guarding for null pointers, or simply by guarding for first == last).

And also fixing it where the issue actually is, and not where it becomes apparent, fixes it also for all new code written against Thrust, where this exact issue would surface.

[dawagner]

Ok, I understand and agree that guarding against first == last deeper in the stack would certainly solve the issue and I appreciate you taking the time of detailing this.

I found some references explaining how accessing std::vector elements may be UB in C++11 (http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_active.html#2182)... Is that what you where thinking about? That's interesting but there's nothing much any of us can do about it.

However I still don't agree on the following:

What is the case against dereferencing a null pointer being UB? The standard says:

null reference cannot exist in a well-defined program, because the only way
to create such a reference would be to bind it to the “object” obtained by indirection through a null pointer,
which causes undefined behavior
(C++14, 8.3.2.5, note)

Regarding memmove: whether C and C++ are good languages or are insane is irrelevant even though I agree. Moreover, this being UB is one aspect of the problem, the other one being that the GNU libc's memmove requires (through an attribute) the source argument to be non-null. That is indeed not necessarily UB but definitely something that prevents using a thrust::vector of initial size 0 safely when using the GNU libc.

[griwes]

I found some references explaining how accessing std::vector elements may be UB in C++11 (http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_active.html#2182)... Is that what you where thinking about? That's interesting but there's nothing much any of us can do about it.

Yes. Also please note that I don't believe the note on the issue is quite correct; the other issue being fixed doesn't fix this.

Regarding null references: there's been a lot of discussion about whether &*first actually dereferences (i.e. forms a reference) or not. I don't think I agree with that position, but I'm not certain the other way either.

And regarding your last paragraph, I think you interpreted what I said wrong (or maybe I wasn't clear enough): I see memmove being specified as UB when the input is a null pointer as the primary problem here, not whether this is through an attribute or through spec. My opinion of this being insane doesn't change the fact I want to fix it ;) Whether we form a null reference or not in the process of calling the function does not really matter, because we invoke UB (even when it is soft, library UB (and by this I mean, the library specification says it is UB, which is slightly different than the core language specifying something as UB, but not in a very significant way)).

[miscco]

I believe this has been fixed.