From 3206dce97fa6aac84d6e46aa07db562f69a93a7b Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 1 May 2024 22:13:25 +0200 Subject: [PATCH] Ensure that {cr}begin works with types that pull in namespace std via ADL Those types must be host only so we can safely SFINAE away our own free functions that would be ambiguous --- .../include/cuda/std/__iterator/access.h | 66 ++++++++++- .../cuda/std/__iterator/reverse_access.h | 55 +++++++-- .../iterator.range/begin-end.pass.cpp | 108 ++++++++++++------ 3 files changed, 180 insertions(+), 49 deletions(-) diff --git a/libcudacxx/include/cuda/std/__iterator/access.h b/libcudacxx/include/cuda/std/__iterator/access.h index 74e4d25d31a..289fd126491 100644 --- a/libcudacxx/include/cuda/std/__iterator/access.h +++ b/libcudacxx/include/cuda/std/__iterator/access.h @@ -21,42 +21,96 @@ # pragma system_header #endif // no system header +#include +#include +#include #include +#include _LIBCUDACXX_BEGIN_NAMESPACE_STD +// We need to detect whether there is already a free function begin / end that would end up being ambiguous. +// This can happen when a type pulls in both namespace std and namespace cuda::std via ADL. +// In that case we are always safe to just not do anything because that type must be host only. +namespace __detect_ambiguous_access +{ +struct __not_ambiguous +{}; +template +_LIBCUDACXX_INLINE_VISIBILITY auto __begin(_Tp (&__array)[_Np]) -> decltype(begin(__array)); +template +_LIBCUDACXX_INLINE_VISIBILITY auto __begin(_Cont& __c) -> decltype(begin(__c)); +template +_LIBCUDACXX_INLINE_VISIBILITY auto __begin(const _Cont& __c) -> decltype(begin(__c)); +_LIBCUDACXX_INLINE_VISIBILITY auto __begin(...) -> __not_ambiguous; +template +struct __begin_not_ambiguous + : is_same())), __not_ambiguous> +{}; +template +struct __begin_not_ambiguous&> : true_type +{}; +template +struct __begin_not_ambiguous&> : true_type +{}; + +template +_LIBCUDACXX_INLINE_VISIBILITY auto __end(_Tp (&__array)[_Np]) -> decltype(end(__array)); +template +_LIBCUDACXX_INLINE_VISIBILITY auto __end(_Cont& __c) -> decltype(end(__c)); +template +_LIBCUDACXX_INLINE_VISIBILITY auto __end(const _Cont& __c) -> decltype(end(__c)); +_LIBCUDACXX_INLINE_VISIBILITY auto __end(...) -> __not_ambiguous; +template +struct __end_not_ambiguous + : is_same())), __not_ambiguous> +{}; +template +struct __end_not_ambiguous&> : true_type +{}; +template +struct __end_not_ambiguous&> : true_type +{}; +} // namespace __detect_ambiguous_access + template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 _Tp* begin(_Tp (&__array)[_Np]) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 auto begin(_Tp (&__array)[_Np]) + -> __enable_if_t<__detect_ambiguous_access::__begin_not_ambiguous<_Tp (&)[_Np]>::value, _Tp*> { return __array; } template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 _Tp* end(_Tp (&__array)[_Np]) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 auto end(_Tp (&__array)[_Np]) + -> __enable_if_t<__detect_ambiguous_access::__end_not_ambiguous<_Tp (&)[_Np]>::value, _Tp*> { return __array + _Np; } template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 auto begin(_Cp& __c) -> decltype(__c.begin()) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 auto begin(_Cp& __c) + -> __enable_if_t<__detect_ambiguous_access::__begin_not_ambiguous<_Cp&>::value, decltype(__c.begin())> { return __c.begin(); } template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 auto begin(const _Cp& __c) -> decltype(__c.begin()) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 auto begin(const _Cp& __c) + -> __enable_if_t<__detect_ambiguous_access::__begin_not_ambiguous::value, decltype(__c.begin())> { return __c.begin(); } template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 auto end(_Cp& __c) -> decltype(__c.end()) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 auto end(_Cp& __c) + -> __enable_if_t<__detect_ambiguous_access::__end_not_ambiguous<_Cp&>::value, decltype(__c.end())> { return __c.end(); } template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 auto end(const _Cp& __c) -> decltype(__c.end()) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 auto end(const _Cp& __c) + -> __enable_if_t<__detect_ambiguous_access::__end_not_ambiguous::value, decltype(__c.end())> { return __c.end(); } diff --git a/libcudacxx/include/cuda/std/__iterator/reverse_access.h b/libcudacxx/include/cuda/std/__iterator/reverse_access.h index dcc9018066e..0df470e9388 100644 --- a/libcudacxx/include/cuda/std/__iterator/reverse_access.h +++ b/libcudacxx/include/cuda/std/__iterator/reverse_access.h @@ -27,16 +27,51 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD -#if _CCCL_STD_VER > 2011 +#if _CCCL_STD_VER >= 2014 + +// We need to detect whether there is already a free function begin / end that would end up being ambiguous. +// This can happen when a type pulls in both namespace std and namespace cuda::std via ADL. +// In that case we are always safe to just not do anything because that type must be host only. +// The exception is initializer list, which must be defined namespace std, so we need to special case that +namespace __detect_ambiguous_raccess +{ +struct __not_ambiguous +{}; +template +_LIBCUDACXX_INLINE_VISIBILITY auto __rbegin(_Tp (&__array)[_Np]) -> decltype(rbegin(__array)); +template +_LIBCUDACXX_INLINE_VISIBILITY auto __rbegin(_Cont& __c) -> decltype(rbegin(__c)); +template +_LIBCUDACXX_INLINE_VISIBILITY auto __rbegin(const _Cont& __c) -> decltype(rbegin(__c)); +_LIBCUDACXX_INLINE_VISIBILITY auto __rbegin(...) -> __not_ambiguous; +template +struct __rbegin_not_ambiguous + : is_same())), __not_ambiguous> +{}; + +template +_LIBCUDACXX_INLINE_VISIBILITY auto __rend(_Tp (&__array)[_Np]) -> decltype(rend(__array)); +template +_LIBCUDACXX_INLINE_VISIBILITY auto __rend(_Cont& __c) -> decltype(rend(__c)); +template +_LIBCUDACXX_INLINE_VISIBILITY auto __rend(const _Cont& __c) -> decltype(rend(__c)); +_LIBCUDACXX_INLINE_VISIBILITY auto __rend(...) -> __not_ambiguous; +template +struct __rend_not_ambiguous + : is_same())), __not_ambiguous> +{}; +} // namespace __detect_ambiguous_raccess template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 reverse_iterator<_Tp*> rbegin(_Tp (&__array)[_Np]) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto rbegin(_Tp (&__array)[_Np]) + -> __enable_if_t<__detect_ambiguous_raccess::__rbegin_not_ambiguous<_Tp (&)[_Np]>::value, reverse_iterator<_Tp*>> { return reverse_iterator<_Tp*>(__array + _Np); } template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 reverse_iterator<_Tp*> rend(_Tp (&__array)[_Np]) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto rend(_Tp (&__array)[_Np]) + -> __enable_if_t<__detect_ambiguous_raccess::__rend_not_ambiguous<_Tp (&)[_Np]>::value, reverse_iterator<_Tp*>> { return reverse_iterator<_Tp*>(__array); } @@ -54,25 +89,29 @@ _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 reverse_iterator } template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto rbegin(_Cp& __c) -> decltype(__c.rbegin()) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto rbegin(_Cp& __c) + -> __enable_if_t<__detect_ambiguous_raccess::__rbegin_not_ambiguous<_Cp&>::value, decltype(__c.rbegin())> { return __c.rbegin(); } template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto rbegin(const _Cp& __c) -> decltype(__c.rbegin()) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto rbegin(const _Cp& __c) + -> __enable_if_t<__detect_ambiguous_raccess::__rbegin_not_ambiguous::value, decltype(__c.rbegin())> { return __c.rbegin(); } template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto rend(_Cp& __c) -> decltype(__c.rend()) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto rend(_Cp& __c) + -> __enable_if_t<__detect_ambiguous_raccess::__rend_not_ambiguous<_Cp&>::value, decltype(__c.rend())> { return __c.rend(); } template -_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto rend(const _Cp& __c) -> decltype(__c.rend()) +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto rend(const _Cp& __c) + -> __enable_if_t<__detect_ambiguous_raccess::__rend_not_ambiguous::value, decltype(__c.rend())> { return __c.rend(); } @@ -89,7 +128,7 @@ _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 auto crend(const _Cp& __c) - return _CUDA_VSTD::rend(__c); } -#endif // _CCCL_STD_VER > 2011 +#endif // _CCCL_STD_VER >= 2014 _LIBCUDACXX_END_NAMESPACE_STD diff --git a/libcudacxx/test/libcudacxx/std/iterators/iterator.range/begin-end.pass.cpp b/libcudacxx/test/libcudacxx/std/iterators/iterator.range/begin-end.pass.cpp index 0daaa35e430..1d9d43b89a3 100644 --- a/libcudacxx/test/libcudacxx/std/iterators/iterator.range/begin-end.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/iterators/iterator.range/begin-end.pass.cpp @@ -41,6 +41,11 @@ #endif #include +#if !defined(TEST_COMPILER_NVRTC) +# include +# include +#endif // !TEST_COMPILER_NVRTC + #include "test_macros.h" // cuda::std::array is explicitly allowed to be initialized with A a = { init-list };. @@ -54,7 +59,7 @@ __host__ __device__ void test_const_container(const C& c, typename C::value_type assert(*cuda::std::begin(c) == val); assert(cuda::std::begin(c) != c.end()); assert(cuda::std::end(c) == c.end()); -#if TEST_STD_VER > 2011 +#if TEST_STD_VER >= 2014 assert(cuda::std::cbegin(c) == c.cbegin()); assert(cuda::std::cbegin(c) != c.cend()); assert(cuda::std::cend(c) == c.cend()); @@ -64,7 +69,7 @@ __host__ __device__ void test_const_container(const C& c, typename C::value_type assert(cuda::std::crbegin(c) == c.crbegin()); assert(cuda::std::crbegin(c) != c.crend()); assert(cuda::std::crend(c) == c.crend()); -#endif +#endif // TEST_STD_VER >= 2014 } template @@ -74,19 +79,11 @@ __host__ __device__ void test_const_container(const cuda::std::initializer_list< assert(*cuda::std::begin(c) == val); assert(cuda::std::begin(c) != c.end()); assert(cuda::std::end(c) == c.end()); -#if TEST_STD_VER > 2011 -// initializer_list doesn't have cbegin/cend/rbegin/rend -// but cuda::std::cbegin(),etc work (b/c they're general fn templates) -// assert ( cuda::std::cbegin(c) == c.cbegin()); -// assert ( cuda::std::cbegin(c) != c.cend()); -// assert ( cuda::std::cend(c) == c.cend()); -// assert ( cuda::std::rbegin(c) == c.rbegin()); -// assert ( cuda::std::rbegin(c) != c.rend()); -// assert ( cuda::std::rend(c) == c.rend()); -// assert ( cuda::std::crbegin(c) == c.crbegin()); -// assert ( cuda::std::crbegin(c) != c.crend()); -// assert ( cuda::std::crend(c) == c.crend()); -#endif +#if TEST_STD_VER >= 2014 + assert(cuda::std::cbegin(c) != cuda::std::cend(c)); + assert(cuda::std::rbegin(c) != cuda::std::rend(c)); + assert(cuda::std::crbegin(c) != cuda::std::crend(c)); +#endif // TEST_STD_VER >= 2014 } template @@ -96,7 +93,7 @@ __host__ __device__ void test_container(C& c, typename C::value_type val) assert(*cuda::std::begin(c) == val); assert(cuda::std::begin(c) != c.end()); assert(cuda::std::end(c) == c.end()); -#if TEST_STD_VER > 2011 +#if TEST_STD_VER >= 2014 assert(cuda::std::cbegin(c) == c.cbegin()); assert(cuda::std::cbegin(c) != c.cend()); assert(cuda::std::cend(c) == c.cend()); @@ -106,7 +103,7 @@ __host__ __device__ void test_container(C& c, typename C::value_type val) assert(cuda::std::crbegin(c) == c.crbegin()); assert(cuda::std::crbegin(c) != c.crend()); assert(cuda::std::crend(c) == c.crend()); -#endif +#endif // TEST_STD_VER >= 2014 } template @@ -116,18 +113,11 @@ __host__ __device__ void test_container(cuda::std::initializer_list& c, T val assert(*cuda::std::begin(c) == val); assert(cuda::std::begin(c) != c.end()); assert(cuda::std::end(c) == c.end()); -#if TEST_STD_VER > 2011 -// initializer_list doesn't have cbegin/cend/rbegin/rend -// assert ( cuda::std::cbegin(c) == c.cbegin()); -// assert ( cuda::std::cbegin(c) != c.cend()); -// assert ( cuda::std::cend(c) == c.cend()); -// assert ( cuda::std::rbegin(c) == c.rbegin()); -// assert ( cuda::std::rbegin(c) != c.rend()); -// assert ( cuda::std::rend(c) == c.rend()); -// assert ( cuda::std::crbegin(c) == c.crbegin()); -// assert ( cuda::std::crbegin(c) != c.crend()); -// assert ( cuda::std::crend(c) == c.crend()); -#endif +#if TEST_STD_VER >= 2014 + assert(cuda::std::cbegin(c) != cuda::std::cend(c)); + assert(cuda::std::rbegin(c) != cuda::std::rend(c)); + assert(cuda::std::crbegin(c) != cuda::std::crend(c)); +#endif // TEST_STD_VER >= 2014 } template @@ -137,12 +127,12 @@ __host__ __device__ void test_const_array(const T (&array)[Sz]) assert(*cuda::std::begin(array) == array[0]); assert(cuda::std::begin(array) != cuda::std::end(array)); assert(cuda::std::end(array) == array + Sz); -#if TEST_STD_VER > 2011 +#if TEST_STD_VER >= 2014 assert(cuda::std::cbegin(array) == array); assert(*cuda::std::cbegin(array) == array[0]); assert(cuda::std::cbegin(array) != cuda::std::cend(array)); assert(cuda::std::cend(array) == array + Sz); -#endif +#endif // TEST_STD_VER >= 2014 } STATIC_TEST_GLOBAL_VAR TEST_CONSTEXPR_GLOBAL int global_array[]{1, 2, 3}; @@ -152,6 +142,52 @@ STATIC_TEST_GLOBAL_VAR TEST_CONSTEXPR_GLOBAL int global_const_array[] = {0, 1, 2 # endif // nvcc > 11.2 #endif // TEST_STD_VER > 2014 +__host__ __device__ void test_ambiguous_std() +{ +#if !defined(TEST_COMPILER_NVRTC) + // clang-format off + NV_IF_TARGET(NV_IS_HOST, ( + { + cuda::std::array<::std::pair, 10> c = {}; + assert(begin(c) == c.begin()); + assert(begin(c) != c.end()); + assert(end(c) == c.end()); + } + + { + cuda::std::initializer_list<::std::pair> init = {{1, 2}}; + assert(begin(init) == init.begin()); + assert(begin(init) != init.end()); + assert(end(init) == init.end()); + } + )) +#if TEST_STD_VER >= 2014 + NV_IF_TARGET(NV_IS_HOST, ( + { + cuda::std::array<::std::pair, 10> c = {}; + assert(cbegin(c) == c.cbegin()); + assert(cbegin(c) != c.cend()); + assert(cend(c) == c.cend()); + assert(rbegin(c) == c.rbegin()); + assert(rbegin(c) != c.rend()); + assert(rend(c) == c.rend()); + assert(crbegin(c) == c.crbegin()); + assert(crbegin(c) != c.crend()); + assert(crend(c) == c.crend()); + } + + { + cuda::std::initializer_list<::std::pair> init = {{1, 2}}; + assert(cbegin(init) != cend(init)); + assert(rbegin(init) != rend(init)); + assert(crbegin(init) != crend(init)); + } + )) +#endif // TEST_STD_VER >= 2014 + // clang-format on +#endif // !TEST_COMPILER_NVRTC +} + int main(int, char**) { #if defined(_LIBCUDACXX_HAS_VECTOR) @@ -185,13 +221,13 @@ int main(int, char**) test_const_container(il, 4); test_const_array(global_array); -#if TEST_STD_VER > 2011 +#if TEST_STD_VER >= 2014 constexpr const int* b = cuda::std::cbegin(global_array); constexpr const int* e = cuda::std::cend(global_array); static_assert(e - b == 3, ""); -#endif +#endif // TEST_STD_VER >= 2014 -#if TEST_STD_VER > 2014 +#if TEST_STD_VER >= 2017 { typedef cuda::std::array C; constexpr const C local_const_array{0, 1, 2, 3, 4}; @@ -228,7 +264,9 @@ int main(int, char**) static_assert(*cuda::std::crbegin(global_const_array) == 4, ""); } # endif // nvcc > 11.2 -#endif // TEST_STD_VER > 2014 +#endif // TEST_STD_VER >= 2017 + + test_ambiguous_std(); return 0; }