Skip to content
This repository has been archived by the owner on Mar 21, 2024. It is now read-only.

Commit

Permalink
Merge pull request #172 from NVIDIA/bugfix/abi_breaking_changes
Browse files Browse the repository at this point in the history
Fix issues in the library that require an ABI break. This introduces ABI version 4.
  • Loading branch information
wmaxey authored Jul 29, 2021
2 parents 2863fde + b302b19 commit 5fe780d
Show file tree
Hide file tree
Showing 7 changed files with 217 additions and 26 deletions.
3 changes: 3 additions & 0 deletions .upstream-tests/test/cuda/pipeline_arrive_on.pass.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down
Original file line number Diff line number Diff line change
@@ -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
//
//===----------------------------------------------------------------------===//

// <cuda/std/complex>

// template<class T>
// class complex
// {
// public:
// typedef T value_type;
// ...
// };

#include <cuda/std/complex>
#include <cuda/std/type_traits>

#include "test_macros.h"

template <class T>
__host__ __device__ void
test()
{
typedef cuda::std::complex<T> C;

static_assert(sizeof(C) == (sizeof(T)*2), "wrong size");
static_assert(alignof(C) == (alignof(T)*2), "misaligned");
}

int main(int, char**)
{
test<float>();
test<double>();
// CUDA treats long double as double
// test<long double>();

return 0;
}
Original file line number Diff line number Diff line change
@@ -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
//
//===----------------------------------------------------------------------===//

// <cuda/std/complex>

// template<class T>
// class complex
// {
// public:
// typedef T value_type;
// ...
// };

#define _LIBCUDACXX_CUDA_ABI_VERSION 3

#include <cuda/std/complex>
#include <cuda/std/type_traits>

#include "test_macros.h"

template <class T>
__host__ __device__ void
test()
{
typedef cuda::std::complex<T> C;

static_assert(sizeof(C) == (sizeof(T)*2), "wrong size");
static_assert(alignof(C) == (alignof(T)), "misaligned");
}

int main(int, char**)
{
test<float>();
test<double>();
// CUDA treats long double as double
// test<long double>();

return 0;
}
Original file line number Diff line number Diff line change
@@ -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
// <cuda/std/chrono>

#pragma nv_diag_suppress declared_but_not_referenced
#pragma nv_diag_suppress set_but_not_used

#define _LIBCUDACXX_CUDA_ABI_VERSION 3

#include <cuda/std/chrono>
#include <cuda/std/type_traits>
#include <cuda/std/cassert>

#include "test_macros.h"
template <typename T>
__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<decltype(3h), cuda::std::chrono::hours>::value, "" );
static_assert ( cuda::std::is_same<decltype(3min), cuda::std::chrono::minutes>::value, "" );
static_assert ( cuda::std::is_same<decltype(3s), cuda::std::chrono::seconds>::value, "" );
static_assert ( cuda::std::is_same<decltype(3ms), cuda::std::chrono::milliseconds>::value, "" );
static_assert ( cuda::std::is_same<decltype(3us), cuda::std::chrono::microseconds>::value, "" );
static_assert ( cuda::std::is_same<decltype(3ns), cuda::std::chrono::nanoseconds>::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<long double, ratio<3600>>::rep >::value, "");
static_assert(cuda::std::is_same< decltype(_min.count()), cuda::std::chrono::duration<long double, ratio< 60>>::rep >::value, "");
// static_assert(cuda::std::is_same< decltype(s.count()), cuda::std::chrono::duration<long double >::rep >::value, "");
static_assert(cuda::std::is_same< decltype(_ms.count()), cuda::std::chrono::duration<long double, milli>::rep >::value, "");
static_assert(cuda::std::is_same< decltype(_us.count()), cuda::std::chrono::duration<long double, micro>::rep >::value, "");
static_assert(cuda::std::is_same< decltype(_ns.count()), cuda::std::chrono::duration<long double, nano>::rep >::value, "");
}

return 0;
}
10 changes: 5 additions & 5 deletions include/cuda/std/detail/__config
Original file line number Diff line number Diff line change
Expand Up @@ -109,21 +109,21 @@
(_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
# define _LIBCUDACXX_CUDA_ABI_VERSION _LIBCUDACXX_CUDA_ABI_VERSION_LATEST
#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
Expand Down
37 changes: 21 additions & 16 deletions libcxx/include/chrono
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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

Expand Down Expand Up @@ -2539,7 +2545,7 @@ inline constexpr days year_month_day::__to_days() const noexcept
static_assert(std::numeric_limits<unsigned>::digits >= 18, "");
static_assert(std::numeric_limits<int>::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<int>(__y) - (__m <= __Feb);
Expand Down Expand Up @@ -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() ?
Expand Down Expand Up @@ -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,
Expand All @@ -3208,9 +3213,9 @@ inline namespace literals
}

_LIBCUDACXX_INLINE_VISIBILITY
constexpr chrono::duration<long double, ratio<3600,1>> operator""h(long double __h)
constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, ratio<3600,1>> operator""h(long double __h)
{
return chrono::duration<long double, ratio<3600,1>>(__h);
return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, ratio<3600,1>>(__h);
}

_LIBCUDACXX_INLINE_VISIBILITY
Expand All @@ -3220,9 +3225,9 @@ inline namespace literals
}

_LIBCUDACXX_INLINE_VISIBILITY
constexpr chrono::duration<long double, ratio<60,1>> operator""min(long double __m)
constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, ratio<60,1>> operator""min(long double __m)
{
return chrono::duration<long double, ratio<60,1>> (__m);
return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, ratio<60,1>> (__m);
}

_LIBCUDACXX_INLINE_VISIBILITY
Expand All @@ -3232,9 +3237,9 @@ inline namespace literals
}

_LIBCUDACXX_INLINE_VISIBILITY
constexpr chrono::duration<long double> operator""s(long double __s)
constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T> operator""s(long double __s)
{
return chrono::duration<long double> (__s);
return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T> (__s);
}

_LIBCUDACXX_INLINE_VISIBILITY
Expand All @@ -3244,9 +3249,9 @@ inline namespace literals
}

_LIBCUDACXX_INLINE_VISIBILITY
constexpr chrono::duration<long double, milli> operator""ms(long double __ms)
constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, milli> operator""ms(long double __ms)
{
return chrono::duration<long double, milli>(__ms);
return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, milli>(__ms);
}

_LIBCUDACXX_INLINE_VISIBILITY
Expand All @@ -3256,9 +3261,9 @@ inline namespace literals
}

_LIBCUDACXX_INLINE_VISIBILITY
constexpr chrono::duration<long double, micro> operator""us(long double __us)
constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, micro> operator""us(long double __us)
{
return chrono::duration<long double, micro> (__us);
return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, micro> (__us);
}

_LIBCUDACXX_INLINE_VISIBILITY
Expand All @@ -3268,9 +3273,9 @@ inline namespace literals
}

_LIBCUDACXX_INLINE_VISIBILITY
constexpr chrono::duration<long double, nano> operator""ns(long double __ns)
constexpr chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, nano> operator""ns(long double __ns)
{
return chrono::duration<long double, nano> (__ns);
return chrono::duration<_LIBCUDACXX_CHRONO_LITERAL_INTERNAL_T, nano> (__ns);
}

#if _LIBCUDACXX_STD_VER > 17 && !defined(_LIBCUDACXX_HAS_NO_CXX20_CHRONO_LITERALS)
Expand Down
16 changes: 11 additions & 5 deletions libcxx/include/complex
Original file line number Diff line number Diff line change
Expand Up @@ -253,9 +253,15 @@ template<class T, class charT, class traits>
#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 _Tp> class _LIBCUDACXX_TEMPLATE_VIS complex;
template<class _Tp> class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_COMPLEX_ALIGNAS(2*sizeof(_Tp)) complex;

template<class _Tp> _LIBCUDACXX_INLINE_VISIBILITY
complex<_Tp> operator*(const complex<_Tp>& __z, const complex<_Tp>& __w);
Expand All @@ -264,7 +270,7 @@ template<class _Tp> _LIBCUDACXX_INLINE_VISIBILITY
complex<_Tp> operator/(const complex<_Tp>& __x, const complex<_Tp>& __y);

template<class _Tp>
class _LIBCUDACXX_TEMPLATE_VIS complex
class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_COMPLEX_ALIGNAS(2*sizeof(_Tp)) complex
{
public:
typedef _Tp value_type;
Expand Down Expand Up @@ -328,7 +334,7 @@ template<> class complex<long double>;
#endif // _LIBCUDACXX_HAS_COMPLEX_LONG_DOUBLE

template<>
class _LIBCUDACXX_TEMPLATE_VIS complex<float>
class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_COMPLEX_ALIGNAS(2*sizeof(float)) complex<float>
{
float __re_;
float __im_;
Expand Down Expand Up @@ -386,7 +392,7 @@ public:
};

template<>
class _LIBCUDACXX_TEMPLATE_VIS complex<double>
class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_COMPLEX_ALIGNAS(2*sizeof(double)) complex<double>
{
double __re_;
double __im_;
Expand Down Expand Up @@ -444,7 +450,7 @@ public:
};

template<>
class _LIBCUDACXX_TEMPLATE_VIS complex<long double>
class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_COMPLEX_ALIGNAS(2*sizeof(long double)) complex<long double>
{
#ifndef _LIBCUDACXX_HAS_COMPLEX_LONG_DOUBLE
public:
Expand Down

0 comments on commit 5fe780d

Please sign in to comment.