Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Try to avoid UB in thrust::reference #2813

Open
wants to merge 2 commits into
base: main
Choose a base branch
from
Open

Try to avoid UB in thrust::reference #2813

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
9763535
Try to avoid UB in `thrust::reference`
miscco Nov 14, 2024
5789899
Make `execute_with_allocator` default constructible
miscco Nov 14, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 4 additions & 5 deletions thrust/testing/memory.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,8 @@
class my_memory_system : public thrust::device_execution_policy<my_memory_system>
{
public:
my_memory_system() = default;

my_memory_system(int)
: correctly_dispatched(false)
, num_copies(0)
Expand All @@ -43,14 +45,11 @@ public:
}

private:
bool correctly_dispatched;
bool correctly_dispatched = true;

// count the number of copies so that we can validate
// that dispatch does not introduce any
unsigned int num_copies;

// disallow default construction
my_memory_system();
unsigned int num_copies = 0;
};

namespace my_old_namespace
Expand Down
5 changes: 5 additions & 0 deletions thrust/thrust/detail/execute_with_allocator_fwd.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,11 @@ struct execute_with_allocator : BaseSystem<execute_with_allocator<Allocator, Bas
Allocator alloc;

public:
_CCCL_HOST_DEVICE execute_with_allocator()
: super_t()
, alloc()
{}

_CCCL_HOST_DEVICE execute_with_allocator(super_t const& super, Allocator alloc_)
: super_t(super)
, alloc(alloc_)
Expand Down
29 changes: 10 additions & 19 deletions thrust/thrust/detail/reference.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -163,10 +163,7 @@ class reference
*/
_CCCL_HOST_DEVICE void swap(derived_type& other)
{
// Avoid default-constructing a system; instead, just use a null pointer
// for dispatch. This assumes that `get_value` will not access any system
// state.
typename thrust::iterator_system<pointer>::type* system = nullptr;
typename thrust::iterator_system<pointer>::type system;
swap(system, other);
}

Expand All @@ -179,10 +176,7 @@ class reference
// about what system the object is on.
_CCCL_HOST_DEVICE operator value_type() const
{
// Avoid default-constructing a system; instead, just use a null pointer
// for dispatch. This assumes that `get_value` will not access any system
// state.
typename thrust::iterator_system<pointer>::type* system = nullptr;
typename thrust::iterator_system<pointer>::type system;
miscco marked this conversation as resolved.
Show resolved Hide resolved
Copy link
Collaborator

@jrhemstad jrhemstad Nov 14, 2024 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can't we just pass the type explicitly instead of materializing a value to pass? That way we can avoid both needing to default construct or using a nullptr.

Suggested change
typename thrust::iterator_system<pointer>::type system;
typename thrust::iterator_system<pointer>::type system;
return convert_to_value_type<decltype(system)>();

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I honestly would also prefer if we did not need to construct systems. From the recent CI failures and the generally open design (users can bring their own systems) we have to assume that users could rely on Thrust not trying to default construct a system. Given the current code "worked" and no data of the system was actually accessed, we can assume that only the type information should be needed. I assume/hope there is a better solution.

Copy link
Contributor

@bernhardmgruber bernhardmgruber Nov 15, 2024 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I wondered how thrust::device_vector produces a system when it inits the elements during construction: It tries to get the system from the allocator. But in case the allocator has no system, it default constructs a system (thrust knows the type of the system for an allocator):

cccl/thrust/thrust/detail/allocator/allocator_traits.inl

Lines 270 to 284 in 87f7246

template <typename Alloc>
_CCCL_HOST_DEVICE ::cuda::std::enable_if_t<has_member_system<Alloc>::value, typename allocator_system<Alloc>::type&>
system(Alloc& a)
{
// return the allocator's system
return a.system();
}
template <typename Alloc>
_CCCL_HOST_DEVICE ::cuda::std::enable_if_t<!has_member_system<Alloc>::value, typename allocator_system<Alloc>::type>
system(Alloc&)
{
// return a copy of a value-initialized system
return typename allocator_system<Alloc>::type();
}

The default allocator for a thrust::device_vector does not seem to contain a system. So it seems we have precedence for default constructing systems when needed.

return convert_to_value_type(system);
}

Expand Down Expand Up @@ -324,10 +318,10 @@ class reference
}

template <typename System>
_CCCL_HOST_DEVICE value_type convert_to_value_type(System* system) const
_CCCL_HOST_DEVICE value_type convert_to_value_type(System& system) const
{
using thrust::system::detail::generic::select_system;
return strip_const_get_value(select_system(*system));
return strip_const_get_value(select_system(system));
}

template <typename System>
Expand All @@ -340,20 +334,17 @@ class reference
}

template <typename System0, typename System1, typename OtherPointer>
_CCCL_HOST_DEVICE void assign_from(System0* system0, System1* system1, OtherPointer src)
_CCCL_HOST_DEVICE void assign_from(System0& system0, System1& system1, OtherPointer src)
{
using thrust::system::detail::generic::select_system;
strip_const_assign_value(select_system(*system0, *system1), src);
strip_const_assign_value(select_system(system0, system1), src);
}

template <typename OtherPointer>
_CCCL_HOST_DEVICE void assign_from(OtherPointer src)
{
// Avoid default-constructing systems; instead, just use a null pointer
// for dispatch. This assumes that `get_value` will not access any system
// state.
typename thrust::iterator_system<pointer>::type* system0 = nullptr;
typename thrust::iterator_system<OtherPointer>::type* system1 = nullptr;
typename thrust::iterator_system<pointer>::type system0;
typename thrust::iterator_system<OtherPointer>::type system1;
assign_from(system0, system1, src);
}

Expand All @@ -367,12 +358,12 @@ class reference
}

template <typename System>
_CCCL_HOST_DEVICE void swap(System* system, derived_type& other)
_CCCL_HOST_DEVICE void swap(System& system, derived_type& other)
{
using thrust::system::detail::generic::iter_swap;
using thrust::system::detail::generic::select_system;

iter_swap(select_system(*system, *system), ptr, other.ptr);
iter_swap(select_system(system, system), ptr, other.ptr);
}
};

Expand Down
Loading