This is an archive of the discontinued LLVM Phabricator instance.

Differential D48036

[CUDA] Make min/max shims host+device.
ClosedPublic

Authored by jlebar on Jun 11 2018, 10:36 AM.

Details

Reviewers
rsmith
Commits
rG5cb41c2acf83: [CUDA] Make min/max shims host+device.
rC336025: [CUDA] Make min/max shims host+device.
rL336025: [CUDA] Make min/max shims host+device.
Summary

Fixes PR37753: min/max can't be called from host device
functions in C++14 mode.

Testcase in a separate test-suite commit.

Diff Detail

Repository
rL LLVM

Event Timeline

jlebar created this revision.Jun 11 2018, 10:36 AM
Herald added a subscriber: sanjoy. · View Herald TranscriptJun 11 2018, 10:36 AM
Harbormaster completed remote builds in B19189: Diff 150790.Jun 11 2018, 10:37 AM
jlebar mentioned this in D48037: [CUDA] Add tests to ensure that std::min/max can be called from __host__ __device__ functions..Jun 11 2018, 10:39 AM
jlebar added a child revision: D48037: [CUDA] Add tests to ensure that std::min/max can be called from __host__ __device__ functions..Jun 11 2018, 10:39 AM
tra added a subscriber: tra.Jun 12 2018, 11:21 AM

Last comment in the bug pointed out that those overloads should be constexpr in c++14. Maybe in a separate patch, though.

https://bugs.llvm.org/show_bug.cgi?id=37753#c5

jlebar added a comment.Jun 13 2018, 9:49 AM

Last comment in the bug pointed out that those overloads should be constexpr in c++14. Maybe in a separate patch, though.

Yeah, would prefer to do it in a separate patch. It's possible that having constexpr min/max in C++14 mode *without a C++14 standard library* will cause problems. (Don't mean to FUD it -- we should try. I just would like to be able to roll them back separately. :)

tra added a comment.Jun 13 2018, 9:59 AM

Ack.

jlebar added a child revision: D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..Jun 13 2018, 2:31 PM
In D48036#1131279, @tra wrote:

Ack.

Patches sent (see dependency chain in phab).

jlebar added a comment.Jun 15 2018, 8:33 AM

@rsmith friendly ping on this one.

jlebar mentioned this in D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..Jun 15 2018, 1:45 PM
jlebar added a comment.Jun 25 2018, 1:08 PM

@rsmith friendly ping on this one.

MaskRay added a subscriber: MaskRay.Jun 29 2018, 3:10 PM
rsmith accepted this revision.Jun 29 2018, 3:15 PM

Looks right to me (other than the missing constexpr in C++14 onwards). Though this is subtle enough that I suspect the only way to know for sure is to try it.

This revision is now accepted and ready to land.Jun 29 2018, 3:15 PM
jlebar added a comment.Jun 29 2018, 3:22 PM

Looks right to me (other than the missing constexpr in C++14 onwards). Though this is subtle enough that I suspect the only way to know for sure is to try it.

Thanks a lot, Richard. FTR the missing constexpr is in D48151.

Closed by commit rL336025: [CUDA] Make min/max shims host+device. (authored by jlebar). · Explain WhyJun 29 2018, 3:32 PM
This revision was automatically updated to reflect the committed changes.
Herald added a subscriber: llvm-commits. · View Herald TranscriptJun 29 2018, 3:32 PM
jlebar mentioned this in rL336029: [CUDA] Add tests to ensure that std::min/max can be called from __host__….Jun 29 2018, 5:37 PM
jlebar mentioned this in rOLDT336029: [CUDA] Add tests to ensure that std::min/max can be called from __host__….Jun 29 2018, 5:42 PM

Revision Contents

PathSize
cfe/
trunk/
lib/
Headers/
cuda_wrappers/
8 lines

Diff 153593

cfe/trunk/lib/Headers/cuda_wrappers/algorithm

Show First 20 Lines • Show All 63 Lines • ▼ Show 20 Lines
namespace std { namespace std {
#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
_GLIBCXX_BEGIN_NAMESPACE_VERSION _GLIBCXX_BEGIN_NAMESPACE_VERSION
#endif #endif
#endif #endif
template <class __T, class __Cmp> template <class __T, class __Cmp>
__attribute__((enable_if(true, ""))) __attribute__((enable_if(true, "")))
inline __device__ const __T & inline __host__ __device__ const __T &
max(const __T &__a, const __T &__b, __Cmp __cmp) { max(const __T &__a, const __T &__b, __Cmp __cmp) {
return __cmp(__a, __b) ? __b : __a; return __cmp(__a, __b) ? __b : __a;
} }
template <class __T> template <class __T>
__attribute__((enable_if(true, ""))) __attribute__((enable_if(true, "")))
inline __device__ const __T & inline __host__ __device__ const __T &
max(const __T &__a, const __T &__b) { max(const __T &__a, const __T &__b) {
return __a < __b ? __b : __a; return __a < __b ? __b : __a;
} }
template <class __T, class __Cmp> template <class __T, class __Cmp>
__attribute__((enable_if(true, ""))) __attribute__((enable_if(true, "")))
inline __device__ const __T & inline __host__ __device__ const __T &
min(const __T &__a, const __T &__b, __Cmp __cmp) { min(const __T &__a, const __T &__b, __Cmp __cmp) {
return __cmp(__b, __a) ? __b : __a; return __cmp(__b, __a) ? __b : __a;
} }
template <class __T> template <class __T>
__attribute__((enable_if(true, ""))) __attribute__((enable_if(true, "")))
inline __device__ const __T & inline __host__ __device__ const __T &
min(const __T &__a, const __T &__b) { min(const __T &__a, const __T &__b) {
return __a < __b ? __a : __b; return __a < __b ? __a : __b;
} }
#ifdef _LIBCPP_END_NAMESPACE_STD #ifdef _LIBCPP_END_NAMESPACE_STD
_LIBCPP_END_NAMESPACE_STD _LIBCPP_END_NAMESPACE_STD
#else #else
#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
_GLIBCXX_END_NAMESPACE_VERSION _GLIBCXX_END_NAMESPACE_VERSION
#endif #endif
} // namespace std } // namespace std
#endif #endif
#endif // __CLANG_CUDA_WRAPPERS_ALGORITHM #endif // __CLANG_CUDA_WRAPPERS_ALGORITHM