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..b809f7284e9 --- /dev/null +++ b/libcudacxx/include/cuda/std/numbers @@ -0,0 +1,162 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#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 + +template +struct __numbers_ill_formed_impl : false_type +{}; + +template +struct __numbers +{ + 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 __numbers<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>> +{ + 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 __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 __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 = __numbers<_Tp>::__e; +template +_CCCL_GLOBAL_CONSTANT _Tp log2e_v = __numbers<_Tp>::__log2e; +template +_CCCL_GLOBAL_CONSTANT _Tp log10e_v = __numbers<_Tp>::__log10e; +template +_CCCL_GLOBAL_CONSTANT _Tp pi_v = __numbers<_Tp>::__pi; +template +_CCCL_GLOBAL_CONSTANT _Tp inv_pi_v = __numbers<_Tp>::__inv_pi; +template +_CCCL_GLOBAL_CONSTANT _Tp inv_sqrtpi_v = __numbers<_Tp>::__inv_sqrtpi; +template +_CCCL_GLOBAL_CONSTANT _Tp ln2_v = __numbers<_Tp>::__ln2; +template +_CCCL_GLOBAL_CONSTANT _Tp ln10_v = __numbers<_Tp>::__ln10; +template +_CCCL_GLOBAL_CONSTANT _Tp sqrt2_v = __numbers<_Tp>::__sqrt2; +template +_CCCL_GLOBAL_CONSTANT _Tp sqrt3_v = __numbers<_Tp>::__sqrt3; +template +_CCCL_GLOBAL_CONSTANT _Tp inv_sqrt3_v = __numbers<_Tp>::__inv_sqrt3; +template +_CCCL_GLOBAL_CONSTANT _Tp egamma_v = __numbers<_Tp>::__egamma; +template +_CCCL_GLOBAL_CONSTANT _Tp phi_v = __numbers<_Tp>::__phi; + +_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; + +} // namespace numbers + +#endif // !_CCCL_NO_VARIABLE_TEMPLATES + +_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/std/numerics/numbers/defined.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/defined.pass.cpp new file mode 100644 index 00000000000..d63d903da34 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/defined.pass.cpp @@ -0,0 +1,86 @@ +//===----------------------------------------------------------------------===// +// +// 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; +} + +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); + 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_type(); + test_type(); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + 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; +} + +int main(int, char**) +{ + test(); + static_assert(test(), ""); + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/numerics/numbers/illformed.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/illformed.fail.cpp new file mode 100644 index 00000000000..0c6e505a27b --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/illformed.fail.cpp @@ -0,0 +1,30 @@ +//===----------------------------------------------------------------------===// +// +// 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:* {{[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}; +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/std/numerics/numbers/user_type.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/user_type.compile.pass.cpp new file mode 100644 index 00000000000..94359f732cd --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/user_type.compile.pass.cpp @@ -0,0 +1,63 @@ +//===----------------------------------------------------------------------===// +// +// 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/std/numerics/numbers/value.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numbers/value.pass.cpp new file mode 100644 index 00000000000..25ba0b9ef71 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/numbers/value.pass.cpp @@ -0,0 +1,131 @@ +//===----------------------------------------------------------------------===// +// +// 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; +} + +// 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; +}