Skip to content

Conversation

@charan-003
Copy link
Contributor

[BUG FIX] Fix nullptr dereference in vector_base::fill_insert for empty vectors

Summary

This PR fixes a nullptr dereference that occurs when calling fill_insert on an empty thrust::vector or thrust::device_vector. The issue was that the function would attempt to access begin(), end(), or position on unallocated storage.

Changes

  1. Added a special case for empty vectors that:
    • Handles allocation if needed
    • Uses uninitialized_fill_n directly to avoid iterator operations on unallocated storage

Fixes #2964

@charan-003 charan-003 requested a review from a team as a code owner September 26, 2025 13:36
@github-project-automation github-project-automation bot moved this to Todo in CCCL Sep 26, 2025
@copy-pr-bot
Copy link
Contributor

copy-pr-bot bot commented Sep 26, 2025

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@cccl-authenticator-app cccl-authenticator-app bot moved this from Todo to In Review in CCCL Sep 26, 2025
@charan-003
Copy link
Contributor Author

Hi @miscco ,

Thank you for your guidance on this issue! I assume that the suggested approach of filling the entire vector with x would change the intended behavior of
fill_insert from "inserting at position" to "resizing and filling.

}

// Handle the simple case: empty vector
if (size() == 0)
Copy link
Contributor

Choose a reason for hiding this comment

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

This is a rare case that should not need a special case, I believe we should only check for capacity < n and then handle all other cases after that

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Hii @miscco

Done! Now clears existing elements and sets size to 0 when n == 0
updated the code to destroy all elements before returning.

@github-project-automation github-project-automation bot moved this from In Review to In Progress in CCCL Sep 29, 2025
@charan-003 charan-003 requested a review from miscco October 1, 2025 15:58
@bernhardmgruber
Copy link
Contributor

/ok to test 08d1ded

@github-actions

This comment has been minimized.

@charan-003
Copy link
Contributor Author

@bernhardmgruber @miscco
Just wondering if there are any changes needed or anything else I should do for this PR?

@bernhardmgruber
Copy link
Contributor

I will let @miscco review this PR next week.

@charan-003
Copy link
Contributor Author

Hi @miscco
Please let me know if any changes are needed for this PR.

@miscco
Copy link
Contributor

miscco commented Nov 3, 2025

/ok to test cb9806f

Comment on lines +885 to +889
if (size() > 0)
{
// we've got room for all of them
// how many existing elements will we displace?
const size_type num_displaced_elements = end() - position;
iterator old_end = end();
m_storage.destroy(begin(), end());
}
m_size = 0;
Copy link
Contributor

Choose a reason for hiding this comment

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

Nitpick:

Suggested change
if (size() > 0)
{
// we've got room for all of them
// how many existing elements will we displace?
const size_type num_displaced_elements = end() - position;
iterator old_end = end();
m_storage.destroy(begin(), end());
}
m_size = 0;
if (size() > 0)
{
m_storage.destroy(begin(), end());
m_size = 0;
}

Comment on lines +900 to +904
// allocate exponentially larger new storage
new_capacity = ::cuda::std::max<size_type>(new_capacity, 2 * capacity());

// finally, fill the range to the insertion point
thrust::fill_n(position, n, x);
} // end if
else
{
// construct new elements at the end of the vector
m_storage.uninitialized_fill_n(end(), n - num_displaced_elements, x);
// do not exceed maximum storage
new_capacity = ::cuda::std::min<size_type>(new_capacity, max_size());
Copy link
Contributor

Choose a reason for hiding this comment

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

I believe we should rather use ::cuda::std::clamp

Suggested change
// allocate exponentially larger new storage
new_capacity = ::cuda::std::max<size_type>(new_capacity, 2 * capacity());
// finally, fill the range to the insertion point
thrust::fill_n(position, n, x);
} // end if
else
{
// construct new elements at the end of the vector
m_storage.uninitialized_fill_n(end(), n - num_displaced_elements, x);
// do not exceed maximum storage
new_capacity = ::cuda::std::min<size_type>(new_capacity, max_size());
// Ensure allocation grows exponentially wiithin bounds
new_capacity = ::cuda::std::clamp<size_type>(new_capacity, static_cast<size_type>(2 * capacity()), max_size());

Comment on lines +906 to +909
if (new_capacity > max_size())
{
throw std::length_error("insert(): insertion exceeds max_size().");
} // end if
Copy link
Contributor

Choose a reason for hiding this comment

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

Important: This cannot happen due to the clamp above

Comment on lines +938 to +941
m_storage.destroy(begin(), end());

storage_type new_storage(copy_allocator_t(), m_storage, new_capacity);
// record the vector's new state
m_storage.swap(new_storage);
Copy link
Contributor

Choose a reason for hiding this comment

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

Nitpick: we should first swap then destroy

Comment on lines +942 to +945
m_size = old_size + n;
}
else
{
Copy link
Contributor

Choose a reason for hiding this comment

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

Nitpick: I would prefer if we would return here and reduce indentation afterwards

Suggested change
m_size = old_size + n;
}
else
{
m_size = old_size + n;
return;
}

m_storage.destroy(new_storage.begin(), new_end);
new_storage.deallocate();
// finally, fill the range to the insertion point
thrust::fill_n(position, n, x);
Copy link
Contributor

Choose a reason for hiding this comment

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

Critical: Missing include

@github-actions
Copy link
Contributor

github-actions bot commented Nov 3, 2025

😬 CI Workflow Results

🟥 Finished in 2h 32m: Pass: 98%/70 | Total: 2d 15h | Max: 2h 31m | Hits: 68%/115298

See results here.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Status: In Progress

Development

Successfully merging this pull request may close these issues.

[BUG] Thrust vector_base::fill_insert does not consider empty initial state

3 participants