Ensure that cuda::std::pair is potentially trivially copyable
#1249
Conversation
trivially copyable is a requirement for memcpy. We want to ensure that our pair implementation satisfies that whenever possible. This is especially important for thrust::pair as that is used in rmm extensively. Fixes NVIDIA#1246
miscco
force-pushed
the
trivial_pair
branch
3 times, most recently
from
814e6b9
to
be1dd81
Compare
|
I am currently investigating why gcc-12 segfaults. I believe this is because we now have a base class which should not have any effect but yeah |
jrhemstad
added
the
backport branch/2.3.x
| !__implicit_default_constructible && _LIBCUDACXX_TRAIT(is_default_constructible, _T1) | ||
| && _LIBCUDACXX_TRAIT(is_default_constructible, _T2); | ||
|
|
||
| static constexpr bool __implicit_copy_constructible_from_T = |
There was a problem hiding this comment.
So at first (just reading the traits) I thought that these 4 should have the _from_T suffix removed from their names, but seeing how they are used, I'd suggest something like __implicit_constructible_from_element_copies and so on; it might be longer, but since there's no T in scope here, even logically, I don't think the current names make much sense. (_from_Ts maybe, but I think that's still muddled in its meaning.)
There was a problem hiding this comment.
How about _from_elements?
There was a problem hiding this comment.
The thing I probably dislike the most about the current name is that it contains "copy_constructible". This has a meaning - this thing we are constructing is created using its copy constructor. But that's not what's happening here, hence my original suggestion.
There was a problem hiding this comment.
Oh yeah, we should just name that __{implicit,explicit}_constructible_from_elements
libcudacxx/include/cuda/std/detail/libcxx/include/__utility/pair.h
Outdated
Show resolved
Hide resolved
libcudacxx/include/cuda/std/detail/libcxx/include/__utility/pair.h
Outdated
Show resolved
Hide resolved
libcudacxx/include/cuda/std/detail/libcxx/include/__utility/pair.h
Outdated
Show resolved
Hide resolved
libcudacxx/include/cuda/std/detail/libcxx/include/__utility/pair.h
Outdated
Show resolved
Hide resolved
libcudacxx/include/cuda/std/detail/libcxx/include/__utility/pair.h
Outdated
Show resolved
Hide resolved
libcudacxx/include/cuda/std/detail/libcxx/include/__utility/pair.h
Outdated
Show resolved
Hide resolved
| // In that case _T1 would be copy assignable but _TestSynthesizeAssignment<_T1> would not | ||
| // This happens e.g for reference types | ||
| template <class _T1> | ||
| struct _TestSynthesizeAssignment |
There was a problem hiding this comment.
question: most types (including helpers) are in lower case. Could you share the motivation behind naming this one in camel calse?
There was a problem hiding this comment.
The thought was that this is nothing that is intended to be ever used outside of that check but yeah that is not really convincing.
I am wondering whether we should just go with _CamelCase for classes and __snake_case for variables?
There was a problem hiding this comment.
Not sure what's the convention in libcu++, in cub it's camel case for template parameters and snake case for everything else (in new style).
There was a problem hiding this comment.
That is what I try to do too, but the issue is that libc++ has a totally different style of everything snake case
I want to be able to see whether something is a type or a function / variable but at the same time I dont want to rewrite the world
libcudacxx/include/cuda/std/detail/libcxx/include/__utility/pair.h
Outdated
Show resolved
Hide resolved
| _T1 first; | ||
| _T2 second; | ||
|
|
||
| template <class _Constraints = __pair_constraints<_T1, _T2>, |
There was a problem hiding this comment.
question: symbols are quite different now:
_ZN4cuda3std3__44pairI1AS3_EC2IS3_S3_NS1_18__pair_constraintsIS3_S3_E15__constructibleIS3_S3_EETnNS1_9enable_ifIXsrT1_24__implicit_constructibleEiE4typeELi0EEEOT_OT0_
vs
_ZN4cuda3std3__44pairI1AS3_EC2IS3_S3_TnPNS1_9enable_ifIXsr10_CheckArgs17__enable_implicitIT_T0_EE5valueEvE4typeELPv0EEEOS7_OS8_
Do we need to bump ABI version?
There was a problem hiding this comment.
I dont think so, we do not change any behavior and worst case it is a link error
There was a problem hiding this comment.
Yeah, I think this is fine. Code that was compiled with the old version will have distinct mangled names from the new version, so they can easily coexist in the same program. As @miscco said, at worst this is a linker error forcing recompilation.
When we do a release, however, this should probably be spelled out as a change potentially requiring recompilation in the release notes.
|
@griwes I have changed the name to |
|
Backport failed for Please cherry-pick the changes locally. git fetch origin branch/2.3.x
git worktree add -d .worktree/backport-1249-to-branch/2.3.x origin/branch/2.3.x
cd .worktree/backport-1249-to-branch/2.3.x
git checkout -b backport-1249-to-branch/2.3.x
ancref=$(git merge-base 50ce545772844ce746134db602af4123cb9cc78e bff2a91eaaf194b59bc3ac7752f13707117f86c5)
git cherry-pick -x $ancref..bff2a91eaaf194b59bc3ac7752f13707117f86c5 |
…IA#1249) trivially copyable is a requirement for memcpy. We want to ensure that our pair implementation satisfies that whenever possible. This is especially important for thrust::pair as that is used in rmm extensively. Fixes NVIDIA#1246 Co-authored-by: Georgy Evtushenko <[email protected]>
…IA#1249) trivially copyable is a requirement for memcpy. We want to ensure that our pair implementation satisfies that whenever possible. This is especially important for thrust::pair as that is used in rmm extensively. Fixes NVIDIA#1246 Co-authored-by: Georgy Evtushenko <[email protected]>
* Ensure that `cuda::std::pair` is potentially trivially copyable (#1249) trivially copyable is a requirement for memcpy. We want to ensure that our pair implementation satisfies that whenever possible. This is especially important for thrust::pair as that is used in rmm extensively. Fixes #1246 Co-authored-by: Georgy Evtushenko <[email protected]> * Mark test as passing --------- Co-authored-by: Georgy Evtushenko <[email protected]>
Patch both CCCL and CUCO to have only internal linkage. For cuco I am working on upstreaming these changes ( NVIDIA/cuCollections#422 ). Once that is accepted and we have validated that moving cuco is stable ( e.g. changes around `cuco::experimental::static_set` ) we can drop this patch set. For cccl the long term fix is to move to CCCL 2.3+, but due to issues ( NVIDIA/cccl#1249, maybe others ) that isn't viable for the 24.02 timeframe. Since the CCCL changes mean C++ and CUDA sources have non compatible ABI's, we need to specify `THRUST_DISABLE_ABI_NAMESPACE` and `THRUST_IGNORE_ABI_NAMESPACE_ERROR` so that we don't change ABI in rapids-cmake consumers since they expect 2.2 behavior. Authors: - Robert Maynard (https://github.com/robertmaynard) Approvers: - Bradley Dice (https://github.com/bdice) URL: #523
miscco
added
backport done
Patch both CCCL and CUCO to have only internal linkage. For cuco I am working on upstreaming these changes ( NVIDIA/cuCollections#422 ). Once that is accepted and we have validated that moving cuco is stable ( e.g. changes around `cuco::experimental::static_set` ) we can drop this patch set. For cccl the long term fix is to move to CCCL 2.3+, but due to issues ( NVIDIA/cccl#1249, maybe others ) that isn't viable for the 24.02 timeframe. Since the CCCL changes mean C++ and CUDA sources have non compatible ABI's, we need to specify `THRUST_DISABLE_ABI_NAMESPACE` and `THRUST_IGNORE_ABI_NAMESPACE_ERROR` so that we don't change ABI in rapids-cmake consumers since they expect 2.2 behavior. Authors: - Robert Maynard (https://github.com/robertmaynard) Approvers: - Bradley Dice (https://github.com/bdice) URL: rapidsai#523
trivially copyable is a requirement for memcpy. We want to ensure that our pair implementation satisfies that whenever possible.
This is especially important for thrust::pair as that is used in rmm extensively.
Fixes #1246