-
Notifications
You must be signed in to change notification settings - Fork 168
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.
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
Replace thrust::swap by cuda::std::swap #2985
base: main
Are you sure you want to change the base?
Conversation
// FIXME(bgruber): swap_allocators already swaps m_allocator, so we are swapping twice here !! | ||
swap_allocators(integral_constant<bool, allocator_traits<Alloc>::propagate_on_container_swap::value>(), | ||
x.m_allocator); | ||
|
||
thrust::swap(m_allocator, x.m_allocator); | ||
// FIXME(bgruber): this should use ADL-two-step swap, but this creates an ambiguity with std::swap until | ||
// https://github.com/NVIDIA/cccl/issues/2984 is resolved. | ||
// swap(m_allocator, x.m_allocator); | ||
::cuda::std::swap(m_allocator, x.m_allocator); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
// FIXME(bgruber): it is really concerning, that swapping an allocator can throw. swap() should be noexcept in | ||
// general. | ||
NV_IF_TARGET(NV_IS_DEVICE, | ||
( | ||
// allocators must be equal when swapping containers with allocators that propagate on swap | ||
assert(!is_allocator_not_equal(other));), | ||
(if (is_allocator_not_equal(other)) { throw allocator_mismatch_on_swap(); })); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I have no ideas about the depths of allocators, but throwing an exception when swapping allocators strikes me as odd.
The reference for std::vector<T>::swap
says:
[...] if
get_allocator() != other.get_allocator()
, the behavior is undefined
I would rather prefer to have a hard fault here, so we can mark all swap functions of thrust vectors as noexcept
.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I mean swapping is move construction + move assignment so it can always potentially throw.
I would say it should not throw from our side, but especially allocators are potentially throwing a ton.
As an example think about a block allocator that reserves a huge slab of memory on construction and hands out slices of it. That could easily fail with an OOM
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I mean swapping is move construction + move assignment so it can always potentially throw.
For a generic implementation, yes. In the long term, relocation could be used to avoid the moved from state, which could simplify things.
As an example think about a block allocator that reserves a huge slab of memory on construction and hands out slices of it. That could easily fail with an OOM
How can swapping two block allocators using memory slabs fail with an OOM? I would still just swap pointers and some bookkeeping?
7a1abe9
to
81287ee
Compare
e9cb9ac
to
6b6050b
Compare
🟨 CI finished in 1h 45m: Pass: 96%/224 | Total: 6d 12h | Avg: 41m 58s | Max: 1h 16m | Hits: 40%/10436
|
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
CUB | |
+/- | Thrust |
CUDA Experimental | |
python | |
CCCL C Parallel Library | |
Catch2Helper |
Modifications in project or dependencies?
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
+/- | Thrust |
CUDA Experimental | |
+/- | python |
+/- | CCCL C Parallel Library |
+/- | Catch2Helper |
🏃 Runner counts (total jobs: 224)
# | Runner |
---|---|
185 | linux-amd64-cpu16 |
16 | linux-arm64-cpu16 |
14 | linux-amd64-gpu-v100-latest-1 |
9 | windows-amd64-cpu16 |
Great, the changes so far break this: thrust::device_reference<T> ref1 = ...;
thrust::device_reference<T> ref2 = ...;
thrust::swap(ref1, ref2); It works with unqualified |
I discussed this with @miscco today and we agreed that the proper way to swap anything in CCCL is: using ::cuda::std::swap;
swap(a, b); This must be made to work for any types of Furthermore, A proper solution requires us to fix: #2984 |
6b6050b
to
db57e74
Compare
52268d2
to
f4566a9
Compare
f4566a9
to
9cb163f
Compare
9cb163f
to
751c208
Compare
Fixes: #2948