-
Notifications
You must be signed in to change notification settings - Fork 186
Fix errors in atomic with small aggregates and enum classes #347
Conversation
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.
Some minor nits about internal helpers. (Allies throughout product code)
include/cuda/std/detail/libcxx/include/support/atomic/atomic_cuda.h
Outdated
Show resolved
Hide resolved
>::type; | ||
|
||
// Arithmetic conversions to/from proxy types | ||
template<class _Tp, _EnableIf<is_arithmetic<_Tp>::value, int> = 0> |
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.
Ditto, we should move to
template<class _Tp, _EnableIf<is_arithmetic<_Tp>::value, int> = 0> | |
template<class _Tp, __enable_if_t<is_arithmetic<_Tp>::value, int> = 0> |
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 was under the impression that the meta base stuff was better to use internally.
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.
That is correct, but my evil plan is to settle on __enable_if_t
globally and then move that over to use the meta based implementation. I really do not like it that we have three different ways to do this.
So anything here is fine by me as that should be a separate PR anyway
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.
That is really strange, do we build in C++03?
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 rebased the branch, which I know is fairly aggressive, but now all the conditionals are using the __blah_t
versions.
|
||
// Arithmetic conversions to/from proxy types | ||
template<class _Tp, _EnableIf<is_arithmetic<_Tp>::value, int> = 0> | ||
constexpr __host__ __device__ inline __cxx_atomic_small_to_32<_Tp> __cxx_small_to_32(_Tp __val) { |
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.
Should this rather use _LIBCUDACXX_INLINE_VISIBILITY
. Also I am not sure if we need the inline here,a s the function is already constexpr
constexpr __host__ __device__ inline __cxx_atomic_small_to_32<_Tp> __cxx_small_to_32(_Tp __val) { | |
_LIBCUDACXX_INLINE_VISIBILITY constexpr __cxx_atomic_small_to_32<_Tp> __cxx_small_to_32(_Tp __val) { |
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.
Updated.
libcudacxx/include/cuda/std/detail/libcxx/include/__type_traits/conditional.h Lines 46 to 54 in 25658c1
Why is it not there? |
Perhaps I just need to update the branch it's based on. I'll do that today. |
Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>
…her than metabase classes
933f51e
to
4a97fc6
Compare
* Add tests and coverage for small enum classes and aggregates * Fix regressions due to improper conversions in small atomic proxies * Use `__conditional_t` alias instead of `conditional` Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com> * /s/__host__ __device__/visibility macro * _[_]conditional_t does not exist, but the metabase _If does * After rebase, use type_trait helpers that provide c++14 backports rather than metabase classes Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>
This PR fixes cases where enum classes and aggregates are assigned to rather than default constructed.
Not that it worked before, but small aggregate types are no longer
constexpr
with this patch.