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

Allow default-initializing Thrust vectors #4183

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

Allow default-initializing Thrust vectors #4183

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
10f28df
Allow default-initializing Thrust vector elements
bernhardmgruber Jul 15, 2024
9e07c71
fix
bernhardmgruber Mar 19, 2025
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
62 changes: 62 additions & 0 deletions thrust/testing/vector.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -834,3 +834,65 @@ void TestVectorMove()
ASSERT_EQUAL(ptr3, ptr4);
}
DECLARE_VECTOR_UNITTEST(TestVectorMove);

struct IntWithInit
{
int value = 42;
};

void TestVectorDefaultInitCtor()
{
// trivially-constructible type: just compilation test, since we cannot check that initialization was skipped
{
thrust::host_vector<int> hv(10, thrust::default_init);
thrust::device_vector<int> dv(10, thrust::default_init);
}

// non-trivially-constructible type: check that initialization was performed
{
thrust::host_vector<IntWithInit> hv(10, thrust::default_init);
for (auto e : hv)
{
ASSERT_EQUAL(e.value, 42);
}

thrust::device_vector<IntWithInit> dv(10, thrust::default_init);
for (auto e : dv)
{
ASSERT_EQUAL(static_cast<IntWithInit>(e).value, 42);
}
}
}
DECLARE_UNITTEST(TestVectorDefaultInitCtor);

void TestVectorDefaultInitResize()
{
// trivially-constructible type: just compilation test, since we cannot check that initialization was skipped
{
thrust::host_vector<int> hv(10);
hv.resize(10, thrust::default_init);
}
{
thrust::device_vector<int> dv(10);
dv.resize(10, thrust::default_init);
}

// non-trivially-constructible type: check that initialization was performed
{
thrust::host_vector<IntWithInit> hv(5);
hv.resize(10, thrust::default_init);
for (auto e : hv)
{
ASSERT_EQUAL(e.value, 42);
}
}
{
thrust::device_vector<IntWithInit> dv(5, thrust::default_init);
dv.resize(10, thrust::default_init);
for (auto e : dv)
{
ASSERT_EQUAL(static_cast<IntWithInit>(e).value, 42);
}
}
}
DECLARE_UNITTEST(TestVectorDefaultInitResize);
24 changes: 23 additions & 1 deletion thrust/thrust/detail/vector_base.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,6 +45,11 @@

THRUST_NAMESPACE_BEGIN

struct default_init_t
Copy link
Contributor

Choose a reason for hiding this comment

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

With cudax we went with uninit_t which I find easier for non-expert users

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I thought default initialization was the safer way, since you cannot skip constructors. I am fine with also providing uninit_t. Would you be fine if I opened a separate PR for that?

Copy link
Contributor

Choose a reason for hiding this comment

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

Honestly I am a bit hesitant here.

The main reason people want to use this is to avoid calling a kernel for instantiation. This breaking for non-trivial types will be a source of frustration / bugs.

We should rather give them the right tool that "works" everywhere, even if it is a sharp edge

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Then I would rather like to provide both.

I see two use cases:

  1. in a generic context, I want to use default_init_t, because I can then write a kernel that just writes to the output buffer by means of assigning to elements, and the assignment is well defined. Buffer creation is faster when the value type is trivially constructible. uninit_t would be a bug here if the value type is not trivially constructible/copyable, since the assignment operator will assign to an object which has not been constructed yet.
  2. for a concrete value type, I want to skip the construction, since I am going to overwrite the data. Depending on the data type, I can just assign to it (trivial) or I need to emplace into the destination.

I wonder though whether the uninit_t should statically assert that the value type is trivially default constructible. Maybe we even need 3 options: default_init_t, uninit_t (checks trivially default constructible), uninit_unchecked_t (skips check, most unsafe and thus more verbose).

{};

inline constexpr default_init_t default_init;

namespace detail
{

Expand Down Expand Up @@ -85,6 +90,11 @@ class vector_base
*/
explicit vector_base(size_type n);

/*! This constructor creates a vector_base with default-initialized elements.
* \param n The number of elements to create.
*/
explicit vector_base(size_type n, default_init_t);

/*! This constructor creates a vector_base with value-initialized elements.
* \param n The number of elements to create.
* \param alloc The allocator to use by this vector_base.
Expand Down Expand Up @@ -208,7 +218,7 @@ class vector_base

/*! \brief Resizes this vector_base to the specified number of elements.
* \param new_size Number of elements this vector_base should contain.
* \throw std::length_error If n exceeds max_size9).
* \throw std::length_error If n exceeds max_size().
*
* This method will resize this vector_base to the specified number of
* elements. If the number is smaller than this vector_base's current
Expand All @@ -217,6 +227,17 @@ class vector_base
*/
void resize(size_type new_size);

/*! \brief Resizes this vector_base to the specified number of elements, performing default-initialization instead of
* value-initialization.
* \param new_size Number of elements this vector_base should contain.
* \throw std::length_error If n exceeds max_size().
*
* This method will resize this vector_base to the specified number of elements. If the number is smaller than this
* vector_base's current size this vector_base is truncated, otherwise this vector_base is extended and new elements
* are default-initialized instead of value-initialized.
*/
void resize(size_type new_size, default_init_t);

/*! \brief Resizes this vector_base to the specified number of elements.
* \param new_size Number of elements this vector_base should contain.
* \param x Data with which new elements should be populated.
Expand Down Expand Up @@ -515,6 +536,7 @@ class vector_base
void insert_dispatch(iterator position, InputIteratorOrIntegralType n, InputIteratorOrIntegralType x, true_type);

// this method appends n value-initialized elements at the end
template <bool OnlyDefaultInit = false>
void append(size_type n);

// this method performs insertion from a fill value
Expand Down
48 changes: 44 additions & 4 deletions thrust/thrust/detail/vector_base.inl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,6 +71,23 @@ vector_base<T, Alloc>::vector_base(size_type n)
value_init(n);
} // end vector_base::vector_base()

template <typename T, typename Alloc>
vector_base<T, Alloc>::vector_base(size_type n, default_init_t)
: m_storage()
, m_size(0)
{
if (n > 0)
{
m_storage.allocate(n);
m_size = n;

if constexpr (!::cuda::std::is_trivially_constructible_v<T>)
{
value_init(n);
}
}
}

template <typename T, typename Alloc>
vector_base<T, Alloc>::vector_base(size_type n, const Alloc& alloc)
: m_storage(alloc)
Expand Down Expand Up @@ -310,6 +327,21 @@ void vector_base<T, Alloc>::resize(size_type new_size)
} // end else
} // end vector_base::resize()

template <typename T, typename Alloc>
void vector_base<T, Alloc>::resize(size_type new_size, default_init_t)
{
if (new_size < size())
{
iterator new_end = begin();
thrust::advance(new_end, new_size);
erase(new_end, end());
} // end if
else
{
append</* OnlyDefaultInit */ true>(new_size - size());
} // end else
} // end vector_base::resize()

template <typename T, typename Alloc>
void vector_base<T, Alloc>::resize(size_type new_size, const value_type& x)
{
Expand Down Expand Up @@ -777,6 +809,7 @@ void vector_base<T, Alloc>::copy_insert(iterator position, ForwardIterator first
} // end vector_base::copy_insert()

template <typename T, typename Alloc>
template <bool OnlyDefaultInit>
void vector_base<T, Alloc>::append(size_type n)
{
if (n != 0)
Expand All @@ -785,8 +818,11 @@ void vector_base<T, Alloc>::append(size_type n)
{
// we've got room for all of them

// default construct new elements at the end of the vector
m_storage.value_initialize_n(end(), n);
if constexpr (!OnlyDefaultInit || !::cuda::std::is_trivially_constructible_v<T>)
{
// default construct new elements at the end of the vector
m_storage.value_initialize_n(end(), n);
}

// extend the size
m_size += n;
Expand Down Expand Up @@ -815,8 +851,12 @@ void vector_base<T, Alloc>::append(size_type n)
// construct copy all elements into the newly allocated storage
new_end = m_storage.uninitialized_copy(begin(), end(), new_storage.begin());

// construct new elements to insert
new_storage.value_initialize_n(new_end, n);
if constexpr (!OnlyDefaultInit || !::cuda::std::is_trivially_constructible_v<T>)
{
// construct new elements to insert
new_storage.value_initialize_n(new_end, n);
}

new_end += n;
} // end try
catch (...)
Expand Down
19 changes: 19 additions & 0 deletions thrust/thrust/device_vector.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -96,6 +96,14 @@ class device_vector : public detail::vector_base<T, Alloc>
: Parent(n)
{}

/*! This constructor creates a \p device_vector with the given size, performing only default-initialization instead of
* value-initialization.
* \param n The number of elements to initially create.
*/
device_vector(size_type n, default_init_t)
: Parent(n, default_init_t{})
{}

/*! This constructor creates a \p device_vector with the given
* size.
* \param n The number of elements to initially create.
Expand Down Expand Up @@ -295,6 +303,17 @@ class device_vector : public detail::vector_base<T, Alloc>
*/
void resize(size_type new_size, const value_type &x = value_type());

/*! \brief Resizes this vector to the specified number of elements, performing default-initialization instead of
* value-initialization.
* \param new_size Number of elements this vector should contain.
* \throw std::length_error If n exceeds max_size().
*
* This method will resize this vector to the specified number of elements. If the number is smaller than this
* vector's current size this vector is truncated, otherwise this vector is extended and new elements are
* default-initialized instead of value-initialized.
*/
void resize(size_type new_size, default_init_t);

/*! Returns the number of elements in this vector.
*/
size_type size() const;
Expand Down
19 changes: 19 additions & 0 deletions thrust/thrust/host_vector.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -97,6 +97,14 @@ class host_vector : public detail::vector_base<T, Alloc>
: Parent(n)
{}

/*! This constructor creates a \p host_vector with the given size, performing only default-initialization instead of
* value-initialization.
* \param n The number of elements to initially create.
*/
_CCCL_HOST host_vector(size_type n, default_init_t)
: Parent(n, default_init_t{})
{}

/*! This constructor creates a \p host_vector with the given
* size.
* \param n The number of elements to initially create.
Expand Down Expand Up @@ -288,6 +296,17 @@ class host_vector : public detail::vector_base<T, Alloc>
*/
void resize(size_type new_size, const value_type &x = value_type());

/*! \brief Resizes this vector to the specified number of elements, performing default-initialization instead of
* value-initialization.
* \param new_size Number of elements this vector should contain.
* \throw std::length_error If n exceeds max_size().
*
* This method will resize this vector to the specified number of elements. If the number is smaller than this
* vector's current size this vector is truncated, otherwise this vector is extended and new elements are
* default-initialized instead of value-initialized.
*/
void resize(size_type new_size, default_init_t);

/*! Returns the number of elements in this vector.
*/
size_type size() const;
Expand Down
Loading