diff --git a/.upstream-tests/test/cuda/pipeline_arrive_on.pass.cpp b/.upstream-tests/test/cuda/pipeline_arrive_on.pass.cpp index 9a9a4e257e..4fab0fdf2b 100644 --- a/.upstream-tests/test/cuda/pipeline_arrive_on.pass.cpp +++ b/.upstream-tests/test/cuda/pipeline_arrive_on.pass.cpp @@ -9,6 +9,9 @@ // UNSUPPORTED: pre-sm-70 +// Remove after bump to version 4 +#define _LIBCUDACXX_CUDA_ABI_VERSION 3 + #pragma nv_diag_suppress static_var_with_dynamic_init #pragma nv_diag_suppress declared_but_not_referenced diff --git a/.upstream-tests/test/std/numerics/complex.number/complex/abi_latest.pass.cpp b/.upstream-tests/test/std/numerics/complex.number/complex/abi_latest.pass.cpp new file mode 100644 index 0000000000..3fe15347e0 --- /dev/null +++ b/.upstream-tests/test/std/numerics/complex.number/complex/abi_latest.pass.cpp @@ -0,0 +1,42 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// + +// template +// class complex +// { +// public: +// typedef T value_type; +// ... +// }; + +#include +#include + +#include "test_macros.h" + +template +__host__ __device__ void +test() +{ + typedef cuda::std::complex C; + + static_assert(sizeof(C) == (sizeof(T)*2), "wrong size"); + static_assert(alignof(C) == (alignof(T)*2), "misaligned"); +} + +int main(int, char**) +{ + test(); + test(); +// CUDA treats long double as double +// test(); + + return 0; +} diff --git a/.upstream-tests/test/std/numerics/complex.number/complex/abi_v3.pass.cpp b/.upstream-tests/test/std/numerics/complex.number/complex/abi_v3.pass.cpp new file mode 100644 index 0000000000..3fcc94604e --- /dev/null +++ b/.upstream-tests/test/std/numerics/complex.number/complex/abi_v3.pass.cpp @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// + +// template +// class complex +// { +// public: +// typedef T value_type; +// ... +// }; + +#define _LIBCUDACXX_CUDA_ABI_VERSION 3 + +#include +#include + +#include "test_macros.h" + +template +__host__ __device__ void +test() +{ + typedef cuda::std::complex C; + + static_assert(sizeof(C) == (sizeof(T)*2), "wrong size"); + static_assert(alignof(C) == (alignof(T)), "misaligned"); +} + +int main(int, char**) +{ + test(); + test(); +// CUDA treats long double as double +// test(); + + return 0; +} diff --git a/.upstream-tests/test/std/utilities/time/time.duration/time.duration.literals/literals.abiv3.pass.cpp b/.upstream-tests/test/std/utilities/time/time.duration/time.duration.literals/literals.abiv3.pass.cpp new file mode 100644 index 0000000000..8dd16c9b65 --- /dev/null +++ b/.upstream-tests/test/std/utilities/time/time.duration/time.duration.literals/literals.abiv3.pass.cpp @@ -0,0 +1,91 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++98, c++03, c++11 +// + +#pragma nv_diag_suppress declared_but_not_referenced +#pragma nv_diag_suppress set_but_not_used + +#define _LIBCUDACXX_CUDA_ABI_VERSION 3 + +#include +#include +#include + +#include "test_macros.h" +template +__host__ __device__ +constexpr bool unused(T &&) {return true;} + +int main(int, char**) +{ + using namespace cuda::std::literals::chrono_literals; + +// long long ABI v3 check + { + constexpr auto _h = 3h; + constexpr auto _min = 3min; + constexpr auto _s = 3s; + constexpr auto _ms = 3ms; + constexpr auto _us = 3us; + constexpr auto _ns = 3ns; + + unused(_h); + unused(_min); + unused(_s); + unused(_ms); + unused(_us); + unused(_ns); + + static_assert(cuda::std::is_same< decltype(_h.count()), cuda::std::chrono::hours::rep >::value, ""); + static_assert(cuda::std::is_same< decltype(_min.count()), cuda::std::chrono::minutes::rep >::value, ""); + static_assert(cuda::std::is_same< decltype(_s.count()), cuda::std::chrono::seconds::rep >::value, ""); + static_assert(cuda::std::is_same< decltype(_ms.count()), cuda::std::chrono::milliseconds::rep >::value, ""); + static_assert(cuda::std::is_same< decltype(_us.count()), cuda::std::chrono::microseconds::rep >::value, ""); + static_assert(cuda::std::is_same< decltype(_ns.count()), cuda::std::chrono::nanoseconds::rep >::value, ""); + + static_assert ( cuda::std::is_same::value, "" ); + static_assert ( cuda::std::is_same::value, "" ); + static_assert ( cuda::std::is_same::value, "" ); + static_assert ( cuda::std::is_same::value, "" ); + static_assert ( cuda::std::is_same::value, "" ); + static_assert ( cuda::std::is_same::value, "" ); + } + +// long double ABI v3 check + { + constexpr auto _h = 3.0h; + constexpr auto _min = 3.0min; + constexpr auto _s = 3.0s; + constexpr auto _ms = 3.0ms; + constexpr auto _us = 3.0us; + constexpr auto _ns = 3.0ns; + + unused(_h); + unused(_min); + unused(_s); + unused(_ms); + unused(_us); + unused(_ns); + + using cuda::std::ratio; + using cuda::std::milli; + using cuda::std::micro; + using cuda::std::nano; + + static_assert(cuda::std::is_same< decltype(_h.count()), cuda::std::chrono::duration>::rep >::value, ""); + static_assert(cuda::std::is_same< decltype(_min.count()), cuda::std::chrono::duration>::rep >::value, ""); + // static_assert(cuda::std::is_same< decltype(s.count()), cuda::std::chrono::duration::rep >::value, ""); + static_assert(cuda::std::is_same< decltype(_ms.count()), cuda::std::chrono::duration::rep >::value, ""); + static_assert(cuda::std::is_same< decltype(_us.count()), cuda::std::chrono::duration::rep >::value, ""); + static_assert(cuda::std::is_same< decltype(_ns.count()), cuda::std::chrono::duration::rep >::value, ""); + } + + return 0; +} diff --git a/include/cuda/std/detail/__config b/include/cuda/std/detail/__config index 5016203cd5..4a2b5a1d1d 100644 --- a/include/cuda/std/detail/__config +++ b/include/cuda/std/detail/__config @@ -109,11 +109,11 @@ (_LIBCUDACXX_CUDA_API_VERSION % 1000) #ifndef _LIBCUDACXX_CUDA_ABI_VERSION_LATEST -# define _LIBCUDACXX_CUDA_ABI_VERSION_LATEST 3 +# define _LIBCUDACXX_CUDA_ABI_VERSION_LATEST 4 #endif #ifdef _LIBCUDACXX_CUDA_ABI_VERSION -# if _LIBCUDACXX_CUDA_ABI_VERSION != 2 && _LIBCUDACXX_CUDA_ABI_VERSION != 3 +# if _LIBCUDACXX_CUDA_ABI_VERSION != 2 && _LIBCUDACXX_CUDA_ABI_VERSION != 3 && _LIBCUDACXX_CUDA_ABI_VERSION != 4 # error Unsupported libcu++ ABI version requested. Please define _LIBCUDACXX_CUDA_ABI_VERSION to either 2 or 3. # endif #else @@ -121,9 +121,9 @@ #endif #ifdef _LIBCUDACXX_PIPELINE_ASSUMED_ABI_VERSION -#if _LIBCUDACXX_PIPELINE_ASSUMED_ABI_VERSION != _LIBCUDACXX_CUDA_ABI_VERSION -#error cuda_pipeline.h has assumed a different libcu++ ABI version than provided by this library. To fix this, please include a libcu++ header before including cuda_pipeline.h, or upgrade to a version of the toolkit this version of libcu++ shipped in. -#endif +# if _LIBCUDACXX_PIPELINE_ASSUMED_ABI_VERSION != _LIBCUDACXX_CUDA_ABI_VERSION +# error cuda_pipeline.h has assumed a different libcu++ ABI version than provided by this library. To fix this, please include a libcu++ header before including cuda_pipeline.h, or upgrade to a version of the toolkit this version of libcu++ shipped in. +# endif #endif #ifndef _LIBCUDACXX_CUDA_ABI_NAMESPACE diff --git a/libcxx/include/chrono b/libcxx/include/chrono index e2939b9536..231a72d978 100644 --- a/libcxx/include/chrono +++ b/libcxx/include/chrono @@ -843,6 +843,12 @@ struct _FilesystemClock; _LIBCUDACXX_END_NAMESPACE_FILESYSTEM #endif // !_LIBCUDACXX_CXX03_LANG +# if _LIBCUDACXX_CUDA_ABI_VERSION > 3 +# define _LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T double +# else +# define _LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T long double +# endif + _LIBCUDACXX_BEGIN_NAMESPACE_STD namespace chrono @@ -1013,7 +1019,7 @@ round(const duration<_Rep, _Period>& __d) return __upper; return __lower.count() & 1 ? __upper : __lower; } -#endif // _LIBCUDACXX_STD_VER > 11 +#endif // _LIBCUDACXX_STD_VER > 11 // duration @@ -2539,7 +2545,7 @@ inline constexpr days year_month_day::__to_days() const noexcept static_assert(std::numeric_limits::digits >= 18, ""); static_assert(std::numeric_limits::digits >= 20 , ""); - // nvcc doesn't allow ODR using constexpr globals. Therefore, + // nvcc doesn't allow ODR using constexpr globals. Therefore, // make a temporary initialized from the global auto constexpr __Feb = February; const int __yr = static_cast(__y) - (__m <= __Feb); @@ -2705,7 +2711,7 @@ chrono::day year_month_day_last::day() const noexcept chrono::day(31), chrono::day(30), chrono::day(31) }; - // nvcc doesn't allow ODR using constexpr globals. Therefore, + // nvcc doesn't allow ODR using constexpr globals. Therefore, // make a temporary initialized from the global auto constexpr __Feb = February; return month() != __Feb || !__y.is_leap() ? @@ -3184,7 +3190,6 @@ constexpr hours make24(const hours& __h, bool __is_pm) noexcept } #endif // _LIBCUDACXX_STD_VER > 11 } // chrono - #if _LIBCUDACXX_STD_VER > 11 // GCC 5 and 6 warn (and then error) on us using the standard reserved UDL names, @@ -3208,9 +3213,9 @@ inline namespace literals } _LIBCUDACXX_INLINE_VISIBILITY - constexpr chrono::duration> operator""h(long double __h) + constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, ratio<3600,1>> operator""h(long double __h) { - return chrono::duration>(__h); + return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, ratio<3600,1>>(__h); } _LIBCUDACXX_INLINE_VISIBILITY @@ -3220,9 +3225,9 @@ inline namespace literals } _LIBCUDACXX_INLINE_VISIBILITY - constexpr chrono::duration> operator""min(long double __m) + constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, ratio<60,1>> operator""min(long double __m) { - return chrono::duration> (__m); + return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, ratio<60,1>> (__m); } _LIBCUDACXX_INLINE_VISIBILITY @@ -3232,9 +3237,9 @@ inline namespace literals } _LIBCUDACXX_INLINE_VISIBILITY - constexpr chrono::duration operator""s(long double __s) + constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T> operator""s(long double __s) { - return chrono::duration (__s); + return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T> (__s); } _LIBCUDACXX_INLINE_VISIBILITY @@ -3244,9 +3249,9 @@ inline namespace literals } _LIBCUDACXX_INLINE_VISIBILITY - constexpr chrono::duration operator""ms(long double __ms) + constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, milli> operator""ms(long double __ms) { - return chrono::duration(__ms); + return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, milli>(__ms); } _LIBCUDACXX_INLINE_VISIBILITY @@ -3256,9 +3261,9 @@ inline namespace literals } _LIBCUDACXX_INLINE_VISIBILITY - constexpr chrono::duration operator""us(long double __us) + constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, micro> operator""us(long double __us) { - return chrono::duration (__us); + return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, micro> (__us); } _LIBCUDACXX_INLINE_VISIBILITY @@ -3268,9 +3273,9 @@ inline namespace literals } _LIBCUDACXX_INLINE_VISIBILITY - constexpr chrono::duration operator""ns(long double __ns) + constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, nano> operator""ns(long double __ns) { - return chrono::duration (__ns); + return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, nano> (__ns); } #if _LIBCUDACXX_STD_VER > 17 && !defined(_LIBCUDACXX_HAS_NO_CXX20_CHRONO_LITERALS) diff --git a/libcxx/include/complex b/libcxx/include/complex index af759fe114..e6a081aaa3 100644 --- a/libcxx/include/complex +++ b/libcxx/include/complex @@ -253,9 +253,15 @@ template #pragma GCC system_header #endif +# if _LIBCUDACXX_CUDA_ABI_VERSION > 3 +# define _LIBCUDACXX_COMPLEX_ALIGNAS(V) _ALIGNAS(V) +# else +# define _LIBCUDACXX_COMPLEX_ALIGNAS(V) +# endif + _LIBCUDACXX_BEGIN_NAMESPACE_STD -template class _LIBCUDACXX_TEMPLATE_VIS complex; +template class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_COMPLEX_ALIGNAS(2*sizeof(_Tp)) complex; template _LIBCUDACXX_INLINE_VISIBILITY complex<_Tp> operator*(const complex<_Tp>& __z, const complex<_Tp>& __w); @@ -264,7 +270,7 @@ template _LIBCUDACXX_INLINE_VISIBILITY complex<_Tp> operator/(const complex<_Tp>& __x, const complex<_Tp>& __y); template -class _LIBCUDACXX_TEMPLATE_VIS complex +class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_COMPLEX_ALIGNAS(2*sizeof(_Tp)) complex { public: typedef _Tp value_type; @@ -328,7 +334,7 @@ template<> class complex; #endif // _LIBCUDACXX_HAS_COMPLEX_LONG_DOUBLE template<> -class _LIBCUDACXX_TEMPLATE_VIS complex +class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_COMPLEX_ALIGNAS(2*sizeof(float)) complex { float __re_; float __im_; @@ -386,7 +392,7 @@ public: }; template<> -class _LIBCUDACXX_TEMPLATE_VIS complex +class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_COMPLEX_ALIGNAS(2*sizeof(double)) complex { double __re_; double __im_; @@ -444,7 +450,7 @@ public: }; template<> -class _LIBCUDACXX_TEMPLATE_VIS complex +class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_COMPLEX_ALIGNAS(2*sizeof(long double)) complex { #ifndef _LIBCUDACXX_HAS_COMPLEX_LONG_DOUBLE public: