From 60485ccbe399426a9bdea0d5d2060f3f21c1d688 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Sat, 11 Jan 2025 11:06:05 +0100 Subject: [PATCH 1/4] implement `` --- docs/libcudacxx/standard_api.rst | 2 + .../standard_api/numerics_library.rst | 4 + .../standard_api/numerics_library/numbers.rst | 9 + libcudacxx/include/cuda/std/numbers | 173 ++++++++++++++++++ libcudacxx/include/cuda/std/version | 4 +- .../aaaaaaaaaaaaa/numbers/defined.pass.cpp | 96 ++++++++++ .../aaaaaaaaaaaaa/numbers/illformed.fail.cpp | 28 +++ .../numbers/user_type.compile.pass.cpp | 61 ++++++ .../aaaaaaaaaaaaa/numbers/value.pass.cpp | 91 +++++++++ 9 files changed, 466 insertions(+), 2 deletions(-) create mode 100644 docs/libcudacxx/standard_api/numerics_library/numbers.rst create mode 100644 libcudacxx/include/cuda/std/numbers create mode 100644 libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/defined.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/illformed.fail.cpp create mode 100644 libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/user_type.compile.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/value.pass.cpp diff --git a/docs/libcudacxx/standard_api.rst b/docs/libcudacxx/standard_api.rst index be806240615..699a815b3ce 100644 --- a/docs/libcudacxx/standard_api.rst +++ b/docs/libcudacxx/standard_api.rst @@ -76,6 +76,8 @@ Feature availability: - C++20 ``std::assume_aligned`` in ```` is available in C++11. +- C++20 ```` are available in C++14. + - C++20 ```` are available in C++17. - all ```` concepts are available in C++17. However, they diff --git a/docs/libcudacxx/standard_api/numerics_library.rst b/docs/libcudacxx/standard_api/numerics_library.rst index 5310cd6ddf9..919e1f5d9c5 100644 --- a/docs/libcudacxx/standard_api/numerics_library.rst +++ b/docs/libcudacxx/standard_api/numerics_library.rst @@ -9,6 +9,7 @@ Numerics Library numerics_library/bit numerics_library/complex + numerics_library/numbers numerics_library/numeric Any Standard C++ header not listed below is omitted. @@ -41,6 +42,9 @@ Any Standard C++ header not listed below is omitted. * - `\ `_ - Fixed-width integer types - libcu++ 1.0.0 / CCCL 2.0.0 / CUDA 10.2 + * - `\ `_ + - Numeric constants + - CCCL 3.0.0 * - `\ `_ - Numeric algorithms - CCCL 2.5.0 diff --git a/docs/libcudacxx/standard_api/numerics_library/numbers.rst b/docs/libcudacxx/standard_api/numerics_library/numbers.rst new file mode 100644 index 00000000000..2a69d3f4d55 --- /dev/null +++ b/docs/libcudacxx/standard_api/numerics_library/numbers.rst @@ -0,0 +1,9 @@ +.. _libcudacxx-standard-api-numerics-numbers: + +```` +====================== + +Extensions +---------- + +- All features of ```` are made available in C++14 onwards diff --git a/libcudacxx/include/cuda/std/numbers b/libcudacxx/include/cuda/std/numbers new file mode 100644 index 00000000000..50e63f86ab5 --- /dev/null +++ b/libcudacxx/include/cuda/std/numbers @@ -0,0 +1,173 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// 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 +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA_STD_NUMBERS +#define _CUDA_STD_NUMBERS + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +namespace numbers +{ + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __e +{ + static constexpr _Tp value = 2.718281828459045235360287471352662; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __log2e +{ + static constexpr _Tp value = 1.442695040888963407359924681001892; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __log10e +{ + static constexpr _Tp value = 0.434294481903251827651128918916605; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __pi +{ + static constexpr _Tp value = 3.141592653589793238462643383279502; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __inv_pi +{ + static constexpr _Tp value = 0.318309886183790671537767526745028; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __inv_sqrtpi +{ + static constexpr _Tp value = 0.564189583547756286948079451560772; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __ln2 +{ + static constexpr _Tp value = 0.693147180559945309417232121458176; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __ln10 +{ + static constexpr _Tp value = 2.302585092994045684017991454684364; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __sqrt2 +{ + static constexpr _Tp value = 1.414213562373095048801688724209698; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __sqrt3 +{ + static constexpr _Tp value = 1.732050807568877293527446341505872; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __inv_sqrt3 +{ + static constexpr _Tp value = 0.577350269189625764509148780501957; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __egamma +{ + static constexpr _Tp value = 0.577215664901532860606512090082402; +}; + +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +struct __phi +{ + static constexpr _Tp value = 1.618033988749894848204586834365638; +}; + +#if !defined(_CCCL_NO_VARIABLE_TEMPLATES) + +template +_CCCL_INLINE_VAR constexpr _Tp e_v = __e<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp log2e_v = __log2e<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp log10e_v = __log10e<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp pi_v = __pi<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp inv_pi_v = __inv_pi<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp inv_sqrtpi_v = __inv_sqrtpi<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp ln2_v = __ln2<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp ln10_v = __ln10<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp sqrt2_v = __sqrt2<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp sqrt3_v = __sqrt3<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp inv_sqrt3_v = __inv_sqrt3<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp egamma_v = __egamma<_Tp>::value; +template +_CCCL_INLINE_VAR constexpr _Tp phi_v = __phi<_Tp>::value; + +_CCCL_INLINE_VAR constexpr double e = e_v; +_CCCL_INLINE_VAR constexpr double log2e = log2e_v; +_CCCL_INLINE_VAR constexpr double log10e = log10e_v; +_CCCL_INLINE_VAR constexpr double pi = pi_v; +_CCCL_INLINE_VAR constexpr double inv_pi = inv_pi_v; +_CCCL_INLINE_VAR constexpr double inv_sqrtpi = inv_sqrtpi_v; +_CCCL_INLINE_VAR constexpr double ln2 = ln2_v; +_CCCL_INLINE_VAR constexpr double ln10 = ln10_v; +_CCCL_INLINE_VAR constexpr double sqrt2 = sqrt2_v; +_CCCL_INLINE_VAR constexpr double sqrt3 = sqrt3_v; +_CCCL_INLINE_VAR constexpr double inv_sqrt3 = inv_sqrt3_v; +_CCCL_INLINE_VAR constexpr double egamma = egamma_v; +_CCCL_INLINE_VAR constexpr double phi = phi_v; + +#endif // !_CCCL_NO_VARIABLE_TEMPLATES + +} // namespace numbers + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _CUDA_STD_NUMBERS diff --git a/libcudacxx/include/cuda/std/version b/libcudacxx/include/cuda/std/version index 951680733c4..b98686e4517 100644 --- a/libcudacxx/include/cuda/std/version +++ b/libcudacxx/include/cuda/std/version @@ -62,7 +62,8 @@ #if !defined(_CCCL_NO_VARIABLE_TEMPLATES) # define __cccl_lib_type_trait_variable_templates 201510L -#endif +# define __cccl_lib_math_constants 201907L +#endif // !_CCCL_NO_VARIABLE_TEMPLATES #if _CCCL_STD_VER >= 2014 # define __cccl_lib_as_const 201510L @@ -205,7 +206,6 @@ // # define __cccl_lib_latch 201907L # endif // !_LIBCUDACXX_HAS_NO_THREADS // # define __cccl_lib_list_remove_return_type 201806L -// # define __cccl_lib_math_constants 201907L // # define __cccl_lib_polymorphic_allocator 201902L // # define __cccl_lib_ranges 201811L // # define __cccl_lib_remove_cvref 201711L diff --git a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/defined.pass.cpp b/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/defined.pass.cpp new file mode 100644 index 00000000000..ea79e812f2f --- /dev/null +++ b/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/defined.pass.cpp @@ -0,0 +1,96 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++11 + +#include + +#include + +template +__host__ __device__ constexpr bool test_defined(const T& value) +{ + ASSERT_SAME_TYPE(ExpectedT, T); + + const ExpectedT* addr = &value; + unused(addr); + + return true; +} + +__host__ __device__ constexpr bool test() +{ + test_defined(cuda::std::numbers::e); + test_defined(cuda::std::numbers::log2e); + test_defined(cuda::std::numbers::log10e); + test_defined(cuda::std::numbers::pi); + test_defined(cuda::std::numbers::inv_pi); + test_defined(cuda::std::numbers::inv_sqrtpi); + test_defined(cuda::std::numbers::ln2); + test_defined(cuda::std::numbers::ln10); + test_defined(cuda::std::numbers::sqrt2); + test_defined(cuda::std::numbers::sqrt3); + test_defined(cuda::std::numbers::inv_sqrt3); + test_defined(cuda::std::numbers::egamma); + test_defined(cuda::std::numbers::phi); + + test_defined(cuda::std::numbers::e_v); + test_defined(cuda::std::numbers::log2e_v); + test_defined(cuda::std::numbers::log10e_v); + test_defined(cuda::std::numbers::pi_v); + test_defined(cuda::std::numbers::inv_pi_v); + test_defined(cuda::std::numbers::inv_sqrtpi_v); + test_defined(cuda::std::numbers::ln2_v); + test_defined(cuda::std::numbers::ln10_v); + test_defined(cuda::std::numbers::sqrt2_v); + test_defined(cuda::std::numbers::sqrt3_v); + test_defined(cuda::std::numbers::inv_sqrt3_v); + test_defined(cuda::std::numbers::egamma_v); + test_defined(cuda::std::numbers::phi_v); + + test_defined(cuda::std::numbers::e_v); + test_defined(cuda::std::numbers::log2e_v); + test_defined(cuda::std::numbers::log10e_v); + test_defined(cuda::std::numbers::pi_v); + test_defined(cuda::std::numbers::inv_pi_v); + test_defined(cuda::std::numbers::inv_sqrtpi_v); + test_defined(cuda::std::numbers::ln2_v); + test_defined(cuda::std::numbers::ln10_v); + test_defined(cuda::std::numbers::sqrt2_v); + test_defined(cuda::std::numbers::sqrt3_v); + test_defined(cuda::std::numbers::inv_sqrt3_v); + test_defined(cuda::std::numbers::egamma_v); + test_defined(cuda::std::numbers::phi_v); + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + test_defined(cuda::std::numbers::e_v); + test_defined(cuda::std::numbers::log2e_v); + test_defined(cuda::std::numbers::log10e_v); + test_defined(cuda::std::numbers::pi_v); + test_defined(cuda::std::numbers::inv_pi_v); + test_defined(cuda::std::numbers::inv_sqrtpi_v); + test_defined(cuda::std::numbers::ln2_v); + test_defined(cuda::std::numbers::ln10_v); + test_defined(cuda::std::numbers::sqrt2_v); + test_defined(cuda::std::numbers::sqrt3_v); + test_defined(cuda::std::numbers::inv_sqrt3_v); + test_defined(cuda::std::numbers::egamma_v); + test_defined(cuda::std::numbers::phi_v); +#endif // !defined(_LIBCUDACXX_NO_LONG_DOUBLE) + + return true; +} + +int main(int, char**) +{ + test(); + static_assert(test(), ""); + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/illformed.fail.cpp b/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/illformed.fail.cpp new file mode 100644 index 00000000000..41180841765 --- /dev/null +++ b/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/illformed.fail.cpp @@ -0,0 +1,28 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++11 + +#include + +// Initializing the primary template is ill-formed. +int log2e{cuda::std::numbers::log2e_v}; // expected-error-re@numbers:* {{static assertion failed{{.*}}A program + // that instantiates a primary template of a mathematical constant variable + // template is ill-formed.}} +int log10e{cuda::std::numbers::log10e_v}; +int pi{cuda::std::numbers::pi_v}; +int inv_pi{cuda::std::numbers::inv_pi_v}; +int inv_sqrtpi{cuda::std::numbers::inv_sqrtpi_v}; +int ln2{cuda::std::numbers::ln2_v}; +int ln10{cuda::std::numbers::ln10_v}; +int sqrt2{cuda::std::numbers::sqrt2_v}; +int sqrt3{cuda::std::numbers::sqrt3_v}; +int inv_sqrt3{cuda::std::numbers::inv_sqrt3_v}; +int egamma{cuda::std::numbers::egamma_v}; +int phi{cuda::std::numbers::phi_v}; diff --git a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/user_type.compile.pass.cpp b/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/user_type.compile.pass.cpp new file mode 100644 index 00000000000..34375450719 --- /dev/null +++ b/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/user_type.compile.pass.cpp @@ -0,0 +1,61 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++11 + +#include + +struct UserType +{ + int value; +}; + +template <> +UserType cuda::std::numbers::e_v{}; + +template <> +UserType cuda::std::numbers::log2e_v{}; + +template <> +UserType cuda::std::numbers::log10e_v{}; + +template <> +UserType cuda::std::numbers::pi_v{}; + +template <> +UserType cuda::std::numbers::inv_pi_v{}; + +template <> +UserType cuda::std::numbers::inv_sqrtpi_v{}; + +template <> +UserType cuda::std::numbers::ln2_v{}; + +template <> +UserType cuda::std::numbers::ln10_v{}; + +template <> +UserType cuda::std::numbers::sqrt2_v{}; + +template <> +UserType cuda::std::numbers::sqrt3_v{}; + +template <> +UserType cuda::std::numbers::inv_sqrt3_v{}; + +template <> +UserType cuda::std::numbers::egamma_v{}; + +template <> +UserType cuda::std::numbers::phi_v{}; + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/value.pass.cpp b/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/value.pass.cpp new file mode 100644 index 00000000000..a183b4b3b88 --- /dev/null +++ b/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/value.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 +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++11 + +#include +#include + +#include + +_CCCL_NV_DIAG_SUPPRESS(1046) // Suppress "floating-point value cannot be represented exactly" + +__host__ __device__ constexpr bool test() +{ + // default constants + assert(cuda::std::numbers::e == 0x1.5bf0a8b145769p+1); + assert(cuda::std::numbers::log2e == 0x1.71547652b82fep+0); + assert(cuda::std::numbers::log10e == 0x1.bcb7b1526e50ep-2); + assert(cuda::std::numbers::pi == 0x1.921fb54442d18p+1); + assert(cuda::std::numbers::inv_pi == 0x1.45f306dc9c883p-2); + assert(cuda::std::numbers::inv_sqrtpi == 0x1.20dd750429b6dp-1); + assert(cuda::std::numbers::ln2 == 0x1.62e42fefa39efp-1); + assert(cuda::std::numbers::ln10 == 0x1.26bb1bbb55516p+1); + assert(cuda::std::numbers::sqrt2 == 0x1.6a09e667f3bcdp+0); + assert(cuda::std::numbers::sqrt3 == 0x1.bb67ae8584caap+0); + assert(cuda::std::numbers::inv_sqrt3 == 0x1.279a74590331cp-1); + assert(cuda::std::numbers::egamma == 0x1.2788cfc6fb619p-1); + assert(cuda::std::numbers::phi == 0x1.9e3779b97f4a8p+0); + + // float constants + assert(cuda::std::numbers::e_v == 0x1.5bf0a8p+1f); + assert(cuda::std::numbers::log2e_v == 0x1.715476p+0f); + assert(cuda::std::numbers::log10e_v == 0x1.bcb7b15p-2f); + assert(cuda::std::numbers::pi_v == 0x1.921fb54p+1f); + assert(cuda::std::numbers::inv_pi_v == 0x1.45f306p-2f); + assert(cuda::std::numbers::inv_sqrtpi_v == 0x1.20dd76p-1f); + assert(cuda::std::numbers::ln2_v == 0x1.62e42fp-1f); + assert(cuda::std::numbers::ln10_v == 0x1.26bb1bp+1f); + assert(cuda::std::numbers::sqrt2_v == 0x1.6a09e6p+0f); + assert(cuda::std::numbers::sqrt3_v == 0x1.bb67aep+0f); + assert(cuda::std::numbers::inv_sqrt3_v == 0x1.279a74p-1f); + assert(cuda::std::numbers::egamma_v == 0x1.2788cfp-1f); + assert(cuda::std::numbers::phi_v == 0x1.9e3779ap+0f); + + // double constants + assert(cuda::std::numbers::e_v == 0x1.5bf0a8b145769p+1); + assert(cuda::std::numbers::log2e_v == 0x1.71547652b82fep+0); + assert(cuda::std::numbers::log10e_v == 0x1.bcb7b1526e50ep-2); + assert(cuda::std::numbers::pi_v == 0x1.921fb54442d18p+1); + assert(cuda::std::numbers::inv_pi_v == 0x1.45f306dc9c883p-2); + assert(cuda::std::numbers::inv_sqrtpi_v == 0x1.20dd750429b6dp-1); + assert(cuda::std::numbers::ln2_v == 0x1.62e42fefa39efp-1); + assert(cuda::std::numbers::ln10_v == 0x1.26bb1bbb55516p+1); + assert(cuda::std::numbers::sqrt2_v == 0x1.6a09e667f3bcdp+0); + assert(cuda::std::numbers::sqrt3_v == 0x1.bb67ae8584caap+0); + assert(cuda::std::numbers::inv_sqrt3_v == 0x1.279a74590331cp-1); + assert(cuda::std::numbers::egamma_v == 0x1.2788cfc6fb619p-1); + assert(cuda::std::numbers::phi_v == 0x1.9e3779b97f4a8p+0); + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + // long double constants + assert(cuda::std::numbers::e_v == 0x1.5bf0a8b145769p+1l); + assert(cuda::std::numbers::log2e_v == 0x1.71547652b82fep+0l); + assert(cuda::std::numbers::log10e_v == 0x1.bcb7b1526e50ep-2l); + assert(cuda::std::numbers::pi_v == 0x1.921fb54442d18p+1l); + assert(cuda::std::numbers::inv_pi_v == 0x1.45f306dc9c883p-2l); + assert(cuda::std::numbers::inv_sqrtpi_v == 0x1.20dd750429b6dp-1l); + assert(cuda::std::numbers::ln2_v == 0x1.62e42fefa39efp-1l); + assert(cuda::std::numbers::ln10_v == 0x1.26bb1bbb55516p+1l); + assert(cuda::std::numbers::sqrt2_v == 0x1.6a09e667f3bcdp+0l); + assert(cuda::std::numbers::sqrt3_v == 0x1.bb67ae8584caap+0l); + assert(cuda::std::numbers::inv_sqrt3_v == 0x1.279a74590331cp-1l); + assert(cuda::std::numbers::egamma_v == 0x1.2788cfc6fb619p-1l); + assert(cuda::std::numbers::phi_v == 0x1.9e3779b97f4a8p+0l); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + + return true; +} + +int main(int, char**) +{ + test(); + static_assert(test(), ""); + return 0; +} From 13546c00311a726be80beca7cae89d4a9ec1c4d7 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Sat, 11 Jan 2025 11:44:40 +0100 Subject: [PATCH 2/4] move test to correct place --- .../{aaaaaaaaaaaaa => std/numerics}/numbers/defined.pass.cpp | 2 ++ .../{aaaaaaaaaaaaa => std/numerics}/numbers/illformed.fail.cpp | 2 ++ .../numerics}/numbers/user_type.compile.pass.cpp | 2 ++ .../{aaaaaaaaaaaaa => std/numerics}/numbers/value.pass.cpp | 2 ++ 4 files changed, 8 insertions(+) rename libcudacxx/test/libcudacxx/{aaaaaaaaaaaaa => std/numerics}/numbers/defined.pass.cpp (99%) rename libcudacxx/test/libcudacxx/{aaaaaaaaaaaaa => std/numerics}/numbers/illformed.fail.cpp (98%) rename libcudacxx/test/libcudacxx/{aaaaaaaaaaaaa => std/numerics}/numbers/user_type.compile.pass.cpp (98%) rename libcudacxx/test/libcudacxx/{aaaaaaaaaaaaa => std/numerics}/numbers/value.pass.cpp (99%) diff --git a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/defined.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/defined.pass.cpp similarity index 99% rename from libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/defined.pass.cpp rename to libcudacxx/test/libcudacxx/std/numerics/numbers/defined.pass.cpp index ea79e812f2f..e8988998316 100644 --- a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/defined.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/defined.pass.cpp @@ -7,6 +7,8 @@ // //===----------------------------------------------------------------------===// +// + // UNSUPPORTED: c++11 #include diff --git a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/illformed.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/illformed.fail.cpp similarity index 98% rename from libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/illformed.fail.cpp rename to libcudacxx/test/libcudacxx/std/numerics/numbers/illformed.fail.cpp index 41180841765..452baac1ba5 100644 --- a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/illformed.fail.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/illformed.fail.cpp @@ -7,6 +7,8 @@ // //===----------------------------------------------------------------------===// +// + // UNSUPPORTED: c++11 #include diff --git a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/user_type.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/user_type.compile.pass.cpp similarity index 98% rename from libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/user_type.compile.pass.cpp rename to libcudacxx/test/libcudacxx/std/numerics/numbers/user_type.compile.pass.cpp index 34375450719..94359f732cd 100644 --- a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/user_type.compile.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/user_type.compile.pass.cpp @@ -7,6 +7,8 @@ // //===----------------------------------------------------------------------===// +// + // UNSUPPORTED: c++11 #include diff --git a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/value.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/value.pass.cpp similarity index 99% rename from libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/value.pass.cpp rename to libcudacxx/test/libcudacxx/std/numerics/numbers/value.pass.cpp index a183b4b3b88..e526250d8e8 100644 --- a/libcudacxx/test/libcudacxx/aaaaaaaaaaaaa/numbers/value.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/value.pass.cpp @@ -7,6 +7,8 @@ // //===----------------------------------------------------------------------===// +// + // UNSUPPORTED: c++11 #include From 22afb87cad9a46ae9d122e693ad0a1a0d3f61a26 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Mon, 13 Jan 2025 08:30:48 +0100 Subject: [PATCH 3/4] add support for extended floating point types --- libcudacxx/include/cuda/std/numbers | 314 +++++++++++++++--- .../std/numerics/numbers/defined.pass.cpp | 70 ++-- .../std/numerics/numbers/value.pass.cpp | 38 +++ 3 files changed, 329 insertions(+), 93 deletions(-) diff --git a/libcudacxx/include/cuda/std/numbers b/libcudacxx/include/cuda/std/numbers index 50e63f86ab5..ee45c3d8748 100644 --- a/libcudacxx/include/cuda/std/numbers +++ b/libcudacxx/include/cuda/std/numbers @@ -22,147 +22,357 @@ #endif // no system header #include +#include #include #include +#if defined(_LIBCUDACXX_HAS_NVFP16) +# include +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_DIAG_PUSH +_CCCL_DIAG_SUPPRESS_CLANG("-Wunused-function") +# include +_CCCL_DIAG_POP +#endif // _LIBCUDACXX_HAS_NVBF16 + _LIBCUDACXX_BEGIN_NAMESPACE_STD namespace numbers { -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __e +{}; + +template +struct __e<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 2.718281828459045235360287471352662; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __log2e +{}; + +template +struct __log2e<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 1.442695040888963407359924681001892; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __log10e +{}; + +template +struct __log10e<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 0.434294481903251827651128918916605; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __pi +{}; + +template +struct __pi<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 3.141592653589793238462643383279502; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __inv_pi +{}; + +template +struct __inv_pi<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 0.318309886183790671537767526745028; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __inv_sqrtpi +{}; + +template +struct __inv_sqrtpi<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 0.564189583547756286948079451560772; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __ln2 +{}; + +template +struct __ln2<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 0.693147180559945309417232121458176; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __ln10 +{}; + +template +struct __ln10<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 2.302585092994045684017991454684364; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __sqrt2 +{}; + +template +struct __sqrt2<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 1.414213562373095048801688724209698; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __sqrt3 +{}; + +template +struct __sqrt3<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 1.732050807568877293527446341505872; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __inv_sqrt3 +{}; + +template +struct __inv_sqrt3<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 0.577350269189625764509148780501957; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __egamma +{}; + +template +struct __egamma<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 0.577215664901532860606512090082402; }; -_CCCL_TEMPLATE(class _Tp) -_CCCL_REQUIRES(_CCCL_TRAIT(is_floating_point, _Tp)) +template struct __phi +{}; + +template +struct __phi<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { static constexpr _Tp value = 1.618033988749894848204586834365638; }; +#if defined(_LIBCUDACXX_HAS_NVFP16) +template <> +struct __e<__half> +{ + static constexpr __half value{__half_raw{0x4170u}}; +}; + +template <> +struct __log2e<__half> +{ + static constexpr __half value{__half_raw{0x3dc5u}}; +}; + +template <> +struct __log10e<__half> +{ + static constexpr __half value{__half_raw{0x36f3u}}; +}; + +template <> +struct __pi<__half> +{ + static constexpr __half value{__half_raw{0x4248u}}; +}; + +template <> +struct __inv_pi<__half> +{ + static constexpr __half value{__half_raw{0x3518u}}; +}; + +template <> +struct __inv_sqrtpi<__half> +{ + static constexpr __half value{__half_raw{0x3883u}}; +}; + +template <> +struct __ln2<__half> +{ + static constexpr __half value{__half_raw{0x398cu}}; +}; + +template <> +struct __ln10<__half> +{ + static constexpr __half value{__half_raw{0x409bu}}; +}; + +template <> +struct __sqrt2<__half> +{ + static constexpr __half value{__half_raw{0x3da8u}}; +}; + +template <> +struct __sqrt3<__half> +{ + static constexpr __half value{__half_raw{0x3eeeu}}; +}; + +template <> +struct __inv_sqrt3<__half> +{ + static constexpr __half value{__half_raw{0x389eu}}; +}; + +template <> +struct __egamma<__half> +{ + static constexpr __half value{__half_raw{0x389eu}}; +}; + +template <> +struct __phi<__half> +{ + static constexpr __half value{__half_raw{0x3e79u}}; +}; +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +template <> +struct __e<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x402eu}}; +}; + +template <> +struct __log2e<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3fb9u}}; +}; + +template <> +struct __log10e<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3edeu}}; +}; + +template <> +struct __pi<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x4049u}}; +}; + +template <> +struct __inv_pi<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3ea3u}}; +}; + +template <> +struct __inv_sqrtpi<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3f10u}}; +}; + +template <> +struct __ln2<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3f31u}}; +}; + +template <> +struct __ln10<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x4013u}}; +}; + +template <> +struct __sqrt2<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3fb5u}}; +}; + +template <> +struct __sqrt3<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3fdeu}}; +}; + +template <> +struct __inv_sqrt3<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3f14u}}; +}; + +template <> +struct __egamma<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3f14u}}; +}; + +template <> +struct __phi<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3fcfu}}; +}; +#endif // _LIBCUDACXX_HAS_NVBF16 + #if !defined(_CCCL_NO_VARIABLE_TEMPLATES) template -_CCCL_INLINE_VAR constexpr _Tp e_v = __e<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp e_v = __e<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp log2e_v = __log2e<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp log2e_v = __log2e<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp log10e_v = __log10e<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp log10e_v = __log10e<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp pi_v = __pi<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp pi_v = __pi<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp inv_pi_v = __inv_pi<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp inv_pi_v = __inv_pi<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp inv_sqrtpi_v = __inv_sqrtpi<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp inv_sqrtpi_v = __inv_sqrtpi<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp ln2_v = __ln2<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp ln2_v = __ln2<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp ln10_v = __ln10<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp ln10_v = __ln10<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp sqrt2_v = __sqrt2<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp sqrt2_v = __sqrt2<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp sqrt3_v = __sqrt3<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp sqrt3_v = __sqrt3<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp inv_sqrt3_v = __inv_sqrt3<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp inv_sqrt3_v = __inv_sqrt3<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp egamma_v = __egamma<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp egamma_v = __egamma<_Tp>::value; template -_CCCL_INLINE_VAR constexpr _Tp phi_v = __phi<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp phi_v = __phi<_Tp>::value; -_CCCL_INLINE_VAR constexpr double e = e_v; -_CCCL_INLINE_VAR constexpr double log2e = log2e_v; -_CCCL_INLINE_VAR constexpr double log10e = log10e_v; -_CCCL_INLINE_VAR constexpr double pi = pi_v; -_CCCL_INLINE_VAR constexpr double inv_pi = inv_pi_v; -_CCCL_INLINE_VAR constexpr double inv_sqrtpi = inv_sqrtpi_v; -_CCCL_INLINE_VAR constexpr double ln2 = ln2_v; -_CCCL_INLINE_VAR constexpr double ln10 = ln10_v; -_CCCL_INLINE_VAR constexpr double sqrt2 = sqrt2_v; -_CCCL_INLINE_VAR constexpr double sqrt3 = sqrt3_v; -_CCCL_INLINE_VAR constexpr double inv_sqrt3 = inv_sqrt3_v; -_CCCL_INLINE_VAR constexpr double egamma = egamma_v; -_CCCL_INLINE_VAR constexpr double phi = phi_v; +_CCCL_GLOBAL_CONSTANT double e = e_v; +_CCCL_GLOBAL_CONSTANT double log2e = log2e_v; +_CCCL_GLOBAL_CONSTANT double log10e = log10e_v; +_CCCL_GLOBAL_CONSTANT double pi = pi_v; +_CCCL_GLOBAL_CONSTANT double inv_pi = inv_pi_v; +_CCCL_GLOBAL_CONSTANT double inv_sqrtpi = inv_sqrtpi_v; +_CCCL_GLOBAL_CONSTANT double ln2 = ln2_v; +_CCCL_GLOBAL_CONSTANT double ln10 = ln10_v; +_CCCL_GLOBAL_CONSTANT double sqrt2 = sqrt2_v; +_CCCL_GLOBAL_CONSTANT double sqrt3 = sqrt3_v; +_CCCL_GLOBAL_CONSTANT double inv_sqrt3 = inv_sqrt3_v; +_CCCL_GLOBAL_CONSTANT double egamma = egamma_v; +_CCCL_GLOBAL_CONSTANT double phi = phi_v; #endif // !_CCCL_NO_VARIABLE_TEMPLATES diff --git a/libcudacxx/test/libcudacxx/std/numerics/numbers/defined.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/defined.pass.cpp index e8988998316..d63d903da34 100644 --- a/libcudacxx/test/libcudacxx/std/numerics/numbers/defined.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/defined.pass.cpp @@ -26,6 +26,26 @@ __host__ __device__ constexpr bool test_defined(const T& value) return true; } +template +__host__ __device__ constexpr bool test_type() +{ + test_defined(cuda::std::numbers::e_v); + test_defined(cuda::std::numbers::log2e_v); + test_defined(cuda::std::numbers::log10e_v); + test_defined(cuda::std::numbers::pi_v); + test_defined(cuda::std::numbers::inv_pi_v); + test_defined(cuda::std::numbers::inv_sqrtpi_v); + test_defined(cuda::std::numbers::ln2_v); + test_defined(cuda::std::numbers::ln10_v); + test_defined(cuda::std::numbers::sqrt2_v); + test_defined(cuda::std::numbers::sqrt3_v); + test_defined(cuda::std::numbers::inv_sqrt3_v); + test_defined(cuda::std::numbers::egamma_v); + test_defined(cuda::std::numbers::phi_v); + + return true; +} + __host__ __device__ constexpr bool test() { test_defined(cuda::std::numbers::e); @@ -42,49 +62,17 @@ __host__ __device__ constexpr bool test() test_defined(cuda::std::numbers::egamma); test_defined(cuda::std::numbers::phi); - test_defined(cuda::std::numbers::e_v); - test_defined(cuda::std::numbers::log2e_v); - test_defined(cuda::std::numbers::log10e_v); - test_defined(cuda::std::numbers::pi_v); - test_defined(cuda::std::numbers::inv_pi_v); - test_defined(cuda::std::numbers::inv_sqrtpi_v); - test_defined(cuda::std::numbers::ln2_v); - test_defined(cuda::std::numbers::ln10_v); - test_defined(cuda::std::numbers::sqrt2_v); - test_defined(cuda::std::numbers::sqrt3_v); - test_defined(cuda::std::numbers::inv_sqrt3_v); - test_defined(cuda::std::numbers::egamma_v); - test_defined(cuda::std::numbers::phi_v); - - test_defined(cuda::std::numbers::e_v); - test_defined(cuda::std::numbers::log2e_v); - test_defined(cuda::std::numbers::log10e_v); - test_defined(cuda::std::numbers::pi_v); - test_defined(cuda::std::numbers::inv_pi_v); - test_defined(cuda::std::numbers::inv_sqrtpi_v); - test_defined(cuda::std::numbers::ln2_v); - test_defined(cuda::std::numbers::ln10_v); - test_defined(cuda::std::numbers::sqrt2_v); - test_defined(cuda::std::numbers::sqrt3_v); - test_defined(cuda::std::numbers::inv_sqrt3_v); - test_defined(cuda::std::numbers::egamma_v); - test_defined(cuda::std::numbers::phi_v); - + test_type(); + test_type(); #if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) - test_defined(cuda::std::numbers::e_v); - test_defined(cuda::std::numbers::log2e_v); - test_defined(cuda::std::numbers::log10e_v); - test_defined(cuda::std::numbers::pi_v); - test_defined(cuda::std::numbers::inv_pi_v); - test_defined(cuda::std::numbers::inv_sqrtpi_v); - test_defined(cuda::std::numbers::ln2_v); - test_defined(cuda::std::numbers::ln10_v); - test_defined(cuda::std::numbers::sqrt2_v); - test_defined(cuda::std::numbers::sqrt3_v); - test_defined(cuda::std::numbers::inv_sqrt3_v); - test_defined(cuda::std::numbers::egamma_v); - test_defined(cuda::std::numbers::phi_v); + test_type(); #endif // !defined(_LIBCUDACXX_NO_LONG_DOUBLE) +#if defined(_LIBCUDACXX_HAS_NVFP16) + test_type<__half>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test_type<__nv_bfloat16>(); +#endif // _LIBCUDACXX_HAS_NVBF16 return true; } diff --git a/libcudacxx/test/libcudacxx/std/numerics/numbers/value.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/value.pass.cpp index e526250d8e8..25ba0b9ef71 100644 --- a/libcudacxx/test/libcudacxx/std/numerics/numbers/value.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/value.pass.cpp @@ -85,9 +85,47 @@ __host__ __device__ constexpr bool test() return true; } +// Extended floating point types are not comparable in constexpr context +__host__ __device__ void test_ext_fp() +{ +#if defined(_LIBCUDACXX_HAS_NVFP16) + // __half constants + assert(cuda::std::numbers::e_v<__half> == __half{2.71875}); + assert(cuda::std::numbers::log2e_v<__half> == __half{1.4423828125}); + assert(cuda::std::numbers::log10e_v<__half> == __half{0.434326171875}); + assert(cuda::std::numbers::pi_v<__half> == __half{3.140625}); + assert(cuda::std::numbers::inv_pi_v<__half> == __half{0.318359375}); + assert(cuda::std::numbers::inv_sqrtpi_v<__half> == __half{0.56396484375}); + assert(cuda::std::numbers::ln2_v<__half> == __half{0.693359375}); + assert(cuda::std::numbers::ln10_v<__half> == __half{2.302734375}); + assert(cuda::std::numbers::sqrt2_v<__half> == __half{1.4140625}); + assert(cuda::std::numbers::sqrt3_v<__half> == __half{1.732421875}); + assert(cuda::std::numbers::inv_sqrt3_v<__half> == __half{0.5771484375}); + assert(cuda::std::numbers::egamma_v<__half> == __half{0.5771484375}); + assert(cuda::std::numbers::phi_v<__half> == __half{1.6181640625}); +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) + assert(cuda::std::numbers::e_v<__nv_bfloat16> == __nv_bfloat16{2.71875f}); + assert(cuda::std::numbers::log2e_v<__nv_bfloat16> == __nv_bfloat16{1.4453125f}); + assert(cuda::std::numbers::log10e_v<__nv_bfloat16> == __nv_bfloat16{0.43359375f}); + assert(cuda::std::numbers::pi_v<__nv_bfloat16> == __nv_bfloat16{3.140625f}); + assert(cuda::std::numbers::inv_pi_v<__nv_bfloat16> == __nv_bfloat16{0.318359375f}); + assert(cuda::std::numbers::inv_sqrtpi_v<__nv_bfloat16> == __nv_bfloat16{0.5625f}); + assert(cuda::std::numbers::ln2_v<__nv_bfloat16> == __nv_bfloat16{0.69140625f}); + assert(cuda::std::numbers::ln10_v<__nv_bfloat16> == __nv_bfloat16{2.296875f}); + assert(cuda::std::numbers::sqrt2_v<__nv_bfloat16> == __nv_bfloat16{1.4140625f}); + assert(cuda::std::numbers::sqrt3_v<__nv_bfloat16> == __nv_bfloat16{1.734375f}); + assert(cuda::std::numbers::inv_sqrt3_v<__nv_bfloat16> == __nv_bfloat16{0.578125f}); + assert(cuda::std::numbers::egamma_v<__nv_bfloat16> == __nv_bfloat16{0.578125f}); + assert(cuda::std::numbers::phi_v<__nv_bfloat16> == __nv_bfloat16{1.6171875f}); +#endif // _LIBCUDACXX_HAS_NVBF16 +} + int main(int, char**) { test(); + test_ext_fp(); static_assert(test(), ""); return 0; } From f211c2fa95f84c08f0484d9d7a90f747609b45be Mon Sep 17 00:00:00 2001 From: David Bayer Date: Mon, 13 Jan 2025 10:36:22 +0100 Subject: [PATCH 4/4] fix review --- libcudacxx/include/cuda/std/numbers | 355 ++++-------------- .../std/numerics/numbers/illformed.fail.cpp | 4 +- 2 files changed, 69 insertions(+), 290 deletions(-) diff --git a/libcudacxx/include/cuda/std/numbers b/libcudacxx/include/cuda/std/numbers index ee45c3d8748..b809f7284e9 100644 --- a/libcudacxx/include/cuda/std/numbers +++ b/libcudacxx/include/cuda/std/numbers @@ -21,7 +21,6 @@ # pragma system_header #endif // no system header -#include #include #include #include @@ -38,327 +37,107 @@ _CCCL_DIAG_POP _LIBCUDACXX_BEGIN_NAMESPACE_STD -namespace numbers -{ - -template -struct __e -{}; - -template -struct __e<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 2.718281828459045235360287471352662; -}; - -template -struct __log2e -{}; - -template -struct __log2e<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 1.442695040888963407359924681001892; -}; - -template -struct __log10e -{}; - -template -struct __log10e<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 0.434294481903251827651128918916605; -}; - -template -struct __pi -{}; - -template -struct __pi<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 3.141592653589793238462643383279502; -}; - -template -struct __inv_pi -{}; - -template -struct __inv_pi<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 0.318309886183790671537767526745028; -}; - -template -struct __inv_sqrtpi -{}; - -template -struct __inv_sqrtpi<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 0.564189583547756286948079451560772; -}; - -template -struct __ln2 -{}; - -template -struct __ln2<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 0.693147180559945309417232121458176; -}; - -template -struct __ln10 -{}; - template -struct __ln10<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 2.302585092994045684017991454684364; -}; - -template -struct __sqrt2 +struct __numbers_ill_formed_impl : false_type {}; -template -struct __sqrt2<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 1.414213562373095048801688724209698; -}; - template -struct __sqrt3 -{}; - -template -struct __sqrt3<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> +struct __numbers { - static constexpr _Tp value = 1.732050807568877293527446341505872; + static_assert(__numbers_ill_formed_impl<_Tp>::value, + "[math.constants] A program that instantiates a primary template of a mathematical constant variable " + "template is ill-formed."); }; -template -struct __inv_sqrt3 -{}; - template -struct __inv_sqrt3<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> +struct __numbers<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> { - static constexpr _Tp value = 0.577350269189625764509148780501957; -}; - -template -struct __egamma -{}; - -template -struct __egamma<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 0.577215664901532860606512090082402; -}; - -template -struct __phi -{}; - -template -struct __phi<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> -{ - static constexpr _Tp value = 1.618033988749894848204586834365638; + static constexpr _Tp __e = 2.718281828459045235360287471352662; + static constexpr _Tp __log2e = 1.442695040888963407359924681001892; + static constexpr _Tp __log10e = 0.434294481903251827651128918916605; + static constexpr _Tp __pi = 3.141592653589793238462643383279502; + static constexpr _Tp __inv_pi = 0.318309886183790671537767526745028; + static constexpr _Tp __inv_sqrtpi = 0.564189583547756286948079451560772; + static constexpr _Tp __ln2 = 0.693147180559945309417232121458176; + static constexpr _Tp __ln10 = 2.302585092994045684017991454684364; + static constexpr _Tp __sqrt2 = 1.414213562373095048801688724209698; + static constexpr _Tp __sqrt3 = 1.732050807568877293527446341505872; + static constexpr _Tp __inv_sqrt3 = 0.577350269189625764509148780501957; + static constexpr _Tp __egamma = 0.577215664901532860606512090082402; + static constexpr _Tp __phi = 1.618033988749894848204586834365638; }; #if defined(_LIBCUDACXX_HAS_NVFP16) template <> -struct __e<__half> -{ - static constexpr __half value{__half_raw{0x4170u}}; -}; - -template <> -struct __log2e<__half> -{ - static constexpr __half value{__half_raw{0x3dc5u}}; -}; - -template <> -struct __log10e<__half> -{ - static constexpr __half value{__half_raw{0x36f3u}}; -}; - -template <> -struct __pi<__half> -{ - static constexpr __half value{__half_raw{0x4248u}}; -}; - -template <> -struct __inv_pi<__half> -{ - static constexpr __half value{__half_raw{0x3518u}}; -}; - -template <> -struct __inv_sqrtpi<__half> -{ - static constexpr __half value{__half_raw{0x3883u}}; -}; - -template <> -struct __ln2<__half> -{ - static constexpr __half value{__half_raw{0x398cu}}; -}; - -template <> -struct __ln10<__half> -{ - static constexpr __half value{__half_raw{0x409bu}}; -}; - -template <> -struct __sqrt2<__half> -{ - static constexpr __half value{__half_raw{0x3da8u}}; -}; - -template <> -struct __sqrt3<__half> -{ - static constexpr __half value{__half_raw{0x3eeeu}}; -}; - -template <> -struct __inv_sqrt3<__half> -{ - static constexpr __half value{__half_raw{0x389eu}}; -}; - -template <> -struct __egamma<__half> -{ - static constexpr __half value{__half_raw{0x389eu}}; -}; - -template <> -struct __phi<__half> -{ - static constexpr __half value{__half_raw{0x3e79u}}; +struct __numbers<__half> +{ + static constexpr __half __e = __half_raw{0x4170u}; + static constexpr __half __log2e = __half_raw{0x3dc5u}; + static constexpr __half __log10e = __half_raw{0x36f3u}; + static constexpr __half __pi = __half_raw{0x4248u}; + static constexpr __half __inv_pi = __half_raw{0x3518u}; + static constexpr __half __inv_sqrtpi = __half_raw{0x3883u}; + static constexpr __half __ln2 = __half_raw{0x398cu}; + static constexpr __half __ln10 = __half_raw{0x409bu}; + static constexpr __half __sqrt2 = __half_raw{0x3da8u}; + static constexpr __half __sqrt3 = __half_raw{0x3eeeu}; + static constexpr __half __inv_sqrt3 = __half_raw{0x389eu}; + static constexpr __half __egamma = __half_raw{0x389eu}; + static constexpr __half __phi = __half_raw{0x3e79u}; }; #endif // _LIBCUDACXX_HAS_NVFP16 #if defined(_LIBCUDACXX_HAS_NVBF16) template <> -struct __e<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x402eu}}; -}; - -template <> -struct __log2e<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3fb9u}}; -}; - -template <> -struct __log10e<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3edeu}}; -}; - -template <> -struct __pi<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x4049u}}; -}; - -template <> -struct __inv_pi<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3ea3u}}; -}; - -template <> -struct __inv_sqrtpi<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3f10u}}; -}; - -template <> -struct __ln2<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3f31u}}; -}; - -template <> -struct __ln10<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x4013u}}; -}; - -template <> -struct __sqrt2<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3fb5u}}; -}; - -template <> -struct __sqrt3<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3fdeu}}; -}; - -template <> -struct __inv_sqrt3<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3f14u}}; -}; - -template <> -struct __egamma<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3f14u}}; -}; - -template <> -struct __phi<__nv_bfloat16> -{ - static constexpr __nv_bfloat16 value{__nv_bfloat16_raw{0x3fcfu}}; +struct __numbers<__nv_bfloat16> +{ + static constexpr __nv_bfloat16 __e = __nv_bfloat16_raw{0x402eu}; + static constexpr __nv_bfloat16 __log2e = __nv_bfloat16_raw{0x3fb9u}; + static constexpr __nv_bfloat16 __log10e = __nv_bfloat16_raw{0x3edeu}; + static constexpr __nv_bfloat16 __pi = __nv_bfloat16_raw{0x4049u}; + static constexpr __nv_bfloat16 __inv_pi = __nv_bfloat16_raw{0x3ea3u}; + static constexpr __nv_bfloat16 __inv_sqrtpi = __nv_bfloat16_raw{0x3f10u}; + static constexpr __nv_bfloat16 __ln2 = __nv_bfloat16_raw{0x3f31u}; + static constexpr __nv_bfloat16 __ln10 = __nv_bfloat16_raw{0x4013u}; + static constexpr __nv_bfloat16 __sqrt2 = __nv_bfloat16_raw{0x3fb5u}; + static constexpr __nv_bfloat16 __sqrt3 = __nv_bfloat16_raw{0x3fdeu}; + static constexpr __nv_bfloat16 __inv_sqrt3 = __nv_bfloat16_raw{0x3f14u}; + static constexpr __nv_bfloat16 __egamma = __nv_bfloat16_raw{0x3f14u}; + static constexpr __nv_bfloat16 __phi = __nv_bfloat16_raw{0x3fcfu}; }; #endif // _LIBCUDACXX_HAS_NVBF16 #if !defined(_CCCL_NO_VARIABLE_TEMPLATES) +namespace numbers +{ + template -_CCCL_GLOBAL_CONSTANT _Tp e_v = __e<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp e_v = __numbers<_Tp>::__e; template -_CCCL_GLOBAL_CONSTANT _Tp log2e_v = __log2e<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp log2e_v = __numbers<_Tp>::__log2e; template -_CCCL_GLOBAL_CONSTANT _Tp log10e_v = __log10e<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp log10e_v = __numbers<_Tp>::__log10e; template -_CCCL_GLOBAL_CONSTANT _Tp pi_v = __pi<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp pi_v = __numbers<_Tp>::__pi; template -_CCCL_GLOBAL_CONSTANT _Tp inv_pi_v = __inv_pi<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp inv_pi_v = __numbers<_Tp>::__inv_pi; template -_CCCL_GLOBAL_CONSTANT _Tp inv_sqrtpi_v = __inv_sqrtpi<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp inv_sqrtpi_v = __numbers<_Tp>::__inv_sqrtpi; template -_CCCL_GLOBAL_CONSTANT _Tp ln2_v = __ln2<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp ln2_v = __numbers<_Tp>::__ln2; template -_CCCL_GLOBAL_CONSTANT _Tp ln10_v = __ln10<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp ln10_v = __numbers<_Tp>::__ln10; template -_CCCL_GLOBAL_CONSTANT _Tp sqrt2_v = __sqrt2<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp sqrt2_v = __numbers<_Tp>::__sqrt2; template -_CCCL_GLOBAL_CONSTANT _Tp sqrt3_v = __sqrt3<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp sqrt3_v = __numbers<_Tp>::__sqrt3; template -_CCCL_GLOBAL_CONSTANT _Tp inv_sqrt3_v = __inv_sqrt3<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp inv_sqrt3_v = __numbers<_Tp>::__inv_sqrt3; template -_CCCL_GLOBAL_CONSTANT _Tp egamma_v = __egamma<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp egamma_v = __numbers<_Tp>::__egamma; template -_CCCL_GLOBAL_CONSTANT _Tp phi_v = __phi<_Tp>::value; +_CCCL_GLOBAL_CONSTANT _Tp phi_v = __numbers<_Tp>::__phi; _CCCL_GLOBAL_CONSTANT double e = e_v; _CCCL_GLOBAL_CONSTANT double log2e = log2e_v; @@ -374,10 +153,10 @@ _CCCL_GLOBAL_CONSTANT double inv_sqrt3 = inv_sqrt3_v; _CCCL_GLOBAL_CONSTANT double egamma = egamma_v; _CCCL_GLOBAL_CONSTANT double phi = phi_v; -#endif // !_CCCL_NO_VARIABLE_TEMPLATES - } // namespace numbers +#endif // !_CCCL_NO_VARIABLE_TEMPLATES + _LIBCUDACXX_END_NAMESPACE_STD #endif // _CUDA_STD_NUMBERS diff --git a/libcudacxx/test/libcudacxx/std/numerics/numbers/illformed.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/illformed.fail.cpp index 452baac1ba5..0c6e505a27b 100644 --- a/libcudacxx/test/libcudacxx/std/numerics/numbers/illformed.fail.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/illformed.fail.cpp @@ -14,8 +14,8 @@ #include // Initializing the primary template is ill-formed. -int log2e{cuda::std::numbers::log2e_v}; // expected-error-re@numbers:* {{static assertion failed{{.*}}A program - // that instantiates a primary template of a mathematical constant variable +int log2e{cuda::std::numbers::log2e_v}; // expected-error-re@numbers:* {{[math.constants] A program that + // instantiates a primary template of a mathematical constant variable // template is ill-formed.}} int log10e{cuda::std::numbers::log10e_v}; int pi{cuda::std::numbers::pi_v};