implement <cuda/std/numbers>

docs/libcudacxx/standard_api.rst

@@ -76,6 +76,8 @@ Feature availability:
 
 -  C++20 ``std::assume_aligned`` in ``<memory>`` is available in C++11.
 
+-  C++20 ``<numbers>`` are available in C++14.
+
 -  C++20 ``<ranges>`` are available in C++17.
 
    -  all ``<ranges>`` concepts are available in C++17. However, they

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.
    * - `\<cuda/std/cstdint\> <https://en.cppreference.com/w/cpp/header/cstdint>`_
      - Fixed-width integer types
      - libcu++ 1.0.0 / CCCL 2.0.0 / CUDA 10.2
+   * - `\<cuda/std/numbers\> <https://en.cppreference.com/w/cpp/header/numbers>`_
+     - Numeric constants
+     - CCCL 3.0.0
    * - `\<cuda/std/numeric\> <https://en.cppreference.com/w/cpp/header/numeric>`_
      - Numeric algorithms
      - CCCL 2.5.0

docs/libcudacxx/standard_api/numerics_library/numbers.rst

@@ -0,0 +1,9 @@
+.. _libcudacxx-standard-api-numerics-numbers:
+
+``<cuda/std/numbers>``
+======================
+
+Extensions
+----------
+
+-  All features of ``<numbers>`` are made available in C++14 onwards

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 <cuda/std/detail/__config>
+
+#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 <cuda/std/__type_traits/enable_if.h>
+#include <cuda/std/__type_traits/is_floating_point.h>
+#include <cuda/std/version>
+
+#if defined(_LIBCUDACXX_HAS_NVFP16)
+#  include <cuda_fp16.h>
+#endif // _LIBCUDACXX_HAS_NVFP16
+#if defined(_LIBCUDACXX_HAS_NVBF16)
+_CCCL_DIAG_PUSH
+_CCCL_DIAG_SUPPRESS_CLANG("-Wunused-function")
+#  include <cuda_bf16.h>
+_CCCL_DIAG_POP
+#endif // _LIBCUDACXX_HAS_NVBF16
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+template <class _Tp>
+struct __numbers_ill_formed_impl : false_type
+{};
+
+template <class _Tp, class = void>
+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 <class _Tp>
+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 <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp e_v = __numbers<_Tp>::__e;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp log2e_v = __numbers<_Tp>::__log2e;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp log10e_v = __numbers<_Tp>::__log10e;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp pi_v = __numbers<_Tp>::__pi;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp inv_pi_v = __numbers<_Tp>::__inv_pi;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp inv_sqrtpi_v = __numbers<_Tp>::__inv_sqrtpi;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp ln2_v = __numbers<_Tp>::__ln2;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp ln10_v = __numbers<_Tp>::__ln10;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp sqrt2_v = __numbers<_Tp>::__sqrt2;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp sqrt3_v = __numbers<_Tp>::__sqrt3;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp inv_sqrt3_v = __numbers<_Tp>::__inv_sqrt3;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp egamma_v = __numbers<_Tp>::__egamma;
+template <class _Tp>
+_CCCL_GLOBAL_CONSTANT _Tp phi_v = __numbers<_Tp>::__phi;
+
+_CCCL_GLOBAL_CONSTANT double e          = e_v<double>;
+_CCCL_GLOBAL_CONSTANT double log2e      = log2e_v<double>;
+_CCCL_GLOBAL_CONSTANT double log10e     = log10e_v<double>;
+_CCCL_GLOBAL_CONSTANT double pi         = pi_v<double>;
+_CCCL_GLOBAL_CONSTANT double inv_pi     = inv_pi_v<double>;
+_CCCL_GLOBAL_CONSTANT double inv_sqrtpi = inv_sqrtpi_v<double>;
+_CCCL_GLOBAL_CONSTANT double ln2        = ln2_v<double>;
+_CCCL_GLOBAL_CONSTANT double ln10       = ln10_v<double>;
+_CCCL_GLOBAL_CONSTANT double sqrt2      = sqrt2_v<double>;
+_CCCL_GLOBAL_CONSTANT double sqrt3      = sqrt3_v<double>;
+_CCCL_GLOBAL_CONSTANT double inv_sqrt3  = inv_sqrt3_v<double>;
+_CCCL_GLOBAL_CONSTANT double egamma     = egamma_v<double>;
+_CCCL_GLOBAL_CONSTANT double phi        = phi_v<double>;
+
+} // namespace numbers
+
+#endif // !_CCCL_NO_VARIABLE_TEMPLATES
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // _CUDA_STD_NUMBERS

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

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.
+//
+//===----------------------------------------------------------------------===//
+
+// <cuda/std/numbers>
+
+// UNSUPPORTED: c++11
+
+#include <cuda/std/numbers>
+
+#include <test_macros.h>
+
+template <class ExpectedT, class T>
+__host__ __device__ constexpr bool test_defined(const T& value)
+{
+  ASSERT_SAME_TYPE(ExpectedT, T);
+
+  const ExpectedT* addr = &value;
+  unused(addr);
+
+  return true;
+}
+
+template <class T>
+__host__ __device__ constexpr bool test_type()
+{
+  test_defined<T>(cuda::std::numbers::e_v<T>);
+  test_defined<T>(cuda::std::numbers::log2e_v<T>);
+  test_defined<T>(cuda::std::numbers::log10e_v<T>);
+  test_defined<T>(cuda::std::numbers::pi_v<T>);
+  test_defined<T>(cuda::std::numbers::inv_pi_v<T>);
+  test_defined<T>(cuda::std::numbers::inv_sqrtpi_v<T>);
+  test_defined<T>(cuda::std::numbers::ln2_v<T>);
+  test_defined<T>(cuda::std::numbers::ln10_v<T>);
+  test_defined<T>(cuda::std::numbers::sqrt2_v<T>);
+  test_defined<T>(cuda::std::numbers::sqrt3_v<T>);
+  test_defined<T>(cuda::std::numbers::inv_sqrt3_v<T>);
+  test_defined<T>(cuda::std::numbers::egamma_v<T>);
+  test_defined<T>(cuda::std::numbers::phi_v<T>);
+
+  return true;
+}
+
+__host__ __device__ constexpr bool test()
+{
+  test_defined<double>(cuda::std::numbers::e);
+  test_defined<double>(cuda::std::numbers::log2e);
+  test_defined<double>(cuda::std::numbers::log10e);
+  test_defined<double>(cuda::std::numbers::pi);
+  test_defined<double>(cuda::std::numbers::inv_pi);
+  test_defined<double>(cuda::std::numbers::inv_sqrtpi);
+  test_defined<double>(cuda::std::numbers::ln2);
+  test_defined<double>(cuda::std::numbers::ln10);
+  test_defined<double>(cuda::std::numbers::sqrt2);
+  test_defined<double>(cuda::std::numbers::sqrt3);
+  test_defined<double>(cuda::std::numbers::inv_sqrt3);
+  test_defined<double>(cuda::std::numbers::egamma);
+  test_defined<double>(cuda::std::numbers::phi);
+
+  test_type<float>();
+  test_type<double>();
+#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE)
+  test_type<long double>();
+#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;
+}

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.
+//
+//===----------------------------------------------------------------------===//
+
+// <cuda/std/numbers>
+
+// UNSUPPORTED: c++11
+
+#include <cuda/std/numbers>
+
+// Initializing the primary template is ill-formed.
+int log2e{cuda::std::numbers::log2e_v<int>}; // 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>};
+int pi{cuda::std::numbers::pi_v<int>};
+int inv_pi{cuda::std::numbers::inv_pi_v<int>};
+int inv_sqrtpi{cuda::std::numbers::inv_sqrtpi_v<int>};
+int ln2{cuda::std::numbers::ln2_v<int>};
+int ln10{cuda::std::numbers::ln10_v<int>};
+int sqrt2{cuda::std::numbers::sqrt2_v<int>};
+int sqrt3{cuda::std::numbers::sqrt3_v<int>};
+int inv_sqrt3{cuda::std::numbers::inv_sqrt3_v<int>};
+int egamma{cuda::std::numbers::egamma_v<int>};
+int phi{cuda::std::numbers::phi_v<int>};

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.
+//
+//===----------------------------------------------------------------------===//
+
+// <cuda/std/numbers>
+
+// UNSUPPORTED: c++11
+
+#include <cuda/std/numbers>
+
+struct UserType
+{
+  int value;
+};
+
+template <>
+UserType cuda::std::numbers::e_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::log2e_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::log10e_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::pi_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::inv_pi_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::inv_sqrtpi_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::ln2_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::ln10_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::sqrt2_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::sqrt3_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::inv_sqrt3_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::egamma_v<UserType>{};
+
+template <>
+UserType cuda::std::numbers::phi_v<UserType>{};
+
+int main(int, char**)
+{
+  return 0;
+}