Skip to content

Commit

Permalink
Try to avoid UB in thrust::reference
Browse files Browse the repository at this point in the history
We were dereferencing a nullptr because we wanted to avoid default constructing a system.

This is obvious UB and prevents us from running the tests with sanitizers, so just default construct the systems.

Addresses #1645
  • Loading branch information
miscco committed Nov 14, 2024
1 parent 100ccde commit ffe0ddb
Show file tree
Hide file tree
Showing 2 changed files with 14 additions and 24 deletions.
9 changes: 4 additions & 5 deletions thrust/testing/memory.cu
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
29 changes: 10 additions & 19 deletions thrust/thrust/detail/reference.h
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{};
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

0 comments on commit ffe0ddb

Please sign in to comment.