Files
derek.holloway 7502bf64a4
CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
Flatten subprojects into a single base monorepo
2026-09-02 16:34:47 -07:00

404 lines
16 KiB
Plaintext

//===----------------------------------------------------------------------===//
//
// 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.
//
//===----------------------------------------------------------------------===//
// MIT License
//
// Modifications Copyright (C) 2025-2026 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in all
// copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.
#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/__floating_point/nvfp_types.h>
#include <cuda/std/__type_traits/enable_if.h>
#include <cuda/std/__type_traits/is_floating_point.h>
#include <cuda/std/version>
#ifdef __HIP_PLATFORM_AMD__
#include <amd/amd_utils.h>
#endif
_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.");
};
_CCCL_DIAG_PUSH
_CCCL_DIAG_SUPPRESS_MSVC(4305) // truncation from 'double' to 'const _Tp'
template <class _Tp>
struct __numbers<_Tp, enable_if_t<_CCCL_TRAIT(is_floating_point, _Tp)>>
{
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __e() noexcept
{
return 2.718281828459045235360287471352662;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __log2e() noexcept
{
return 1.442695040888963407359924681001892;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __log10e() noexcept
{
return 0.434294481903251827651128918916605;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __pi() noexcept
{
return 3.141592653589793238462643383279502;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __inv_pi() noexcept
{
return 0.318309886183790671537767526745028;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __inv_sqrtpi() noexcept
{
return 0.564189583547756286948079451560772;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __ln2() noexcept
{
return 0.693147180559945309417232121458176;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __ln10() noexcept
{
return 2.302585092994045684017991454684364;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __sqrt2() noexcept
{
return 1.414213562373095048801688724209698;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __sqrt3() noexcept
{
return 1.732050807568877293527446341505872;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __inv_sqrt3() noexcept
{
return 0.577350269189625764509148780501957;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __egamma() noexcept
{
return 0.577215664901532860606512090082402;
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __phi() noexcept
{
return 1.618033988749894848204586834365638;
}
};
_CCCL_DIAG_POP
#if _LIBCUDACXX_HAS_NVFP16()
// NOTE(HIP/AMD): for ROCm 7.10 and earlier constexpression setting of __half values is not possible
# if LIBHIPCXX_ROCM_VERSION_GE(7, 11, 0)
template <>
struct __numbers<__half>
{
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __e() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x4170u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __log2e() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x3dc5u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __log10e() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x36f3u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __pi() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x4248u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __inv_pi() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x3518u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __inv_sqrtpi() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x3883u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __ln2() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x398cu};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __ln10() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x409bu};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __sqrt2() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x3da8u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __sqrt3() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x3eeeu};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __inv_sqrt3() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x389eu};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __egamma() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x389eu};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __half __phi() noexcept
{
// TODO(HIP/AMD): remove this WAR once __half_raw has only the x member
// This WAR is necessary as otherwise the data member will be active and the
// passed value will be interpreted as a FLOAT16 instead of an unsigned int
return __half_raw{.x = 0x3e79u};
}
};
# endif // LIBHIPCXX_ROCM_VERSION_GE(7, 11, 0)
#endif // _LIBCUDACXX_HAS_NVFP16()
#if _LIBCUDACXX_HAS_NVBF16()
template <>
struct __numbers<__hip_bfloat16>
{
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __e() noexcept
{
return __hip_bfloat16_raw{0x402eu};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __log2e() noexcept
{
return __hip_bfloat16_raw{0x3fb9u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __log10e() noexcept
{
return __hip_bfloat16_raw{0x3edeu};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __pi() noexcept
{
return __hip_bfloat16_raw{0x4049u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __inv_pi() noexcept
{
return __hip_bfloat16_raw{0x3ea3u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __inv_sqrtpi() noexcept
{
return __hip_bfloat16_raw{0x3f10u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __ln2() noexcept
{
return __hip_bfloat16_raw{0x3f31u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __ln10() noexcept
{
return __hip_bfloat16_raw{0x4013u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __sqrt2() noexcept
{
return __hip_bfloat16_raw{0x3fb5u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __sqrt3() noexcept
{
return __hip_bfloat16_raw{0x3fdeu};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __inv_sqrt3() noexcept
{
return __hip_bfloat16_raw{0x3f14u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __egamma() noexcept
{
return __hip_bfloat16_raw{0x3f14u};
}
static _LIBCUDACXX_HIDE_FROM_ABI constexpr __hip_bfloat16 __phi() noexcept
{
return __hip_bfloat16_raw{0x3fcfu};
}
};
#endif // _LIBCUDACXX_HAS_NVBF16()
namespace numbers
{
template <class _Tp>
inline constexpr _Tp e_v = __numbers<_Tp>::__e();
template <class _Tp>
inline constexpr _Tp log2e_v = __numbers<_Tp>::__log2e();
template <class _Tp>
inline constexpr _Tp log10e_v = __numbers<_Tp>::__log10e();
template <class _Tp>
inline constexpr _Tp pi_v = __numbers<_Tp>::__pi();
template <class _Tp>
inline constexpr _Tp inv_pi_v = __numbers<_Tp>::__inv_pi();
template <class _Tp>
inline constexpr _Tp inv_sqrtpi_v = __numbers<_Tp>::__inv_sqrtpi();
template <class _Tp>
inline constexpr _Tp ln2_v = __numbers<_Tp>::__ln2();
template <class _Tp>
inline constexpr _Tp ln10_v = __numbers<_Tp>::__ln10();
template <class _Tp>
inline constexpr _Tp sqrt2_v = __numbers<_Tp>::__sqrt2();
template <class _Tp>
inline constexpr _Tp sqrt3_v = __numbers<_Tp>::__sqrt3();
template <class _Tp>
inline constexpr _Tp inv_sqrt3_v = __numbers<_Tp>::__inv_sqrt3();
template <class _Tp>
inline constexpr _Tp egamma_v = __numbers<_Tp>::__egamma();
template <class _Tp>
inline constexpr _Tp phi_v = __numbers<_Tp>::__phi();
#if !_CCCL_COMPILER(MSVC)
// MSVC errors here because of "error: A __device__ variable template cannot have a const qualified type on Windows"
# if _LIBCUDACXX_HAS_NVFP16()
// NOTE(HIP/AMD): for ROCm 7.10 and earlier constexpression setting of __half values is not possible
# if LIBHIPCXX_ROCM_VERSION_GE(7, 11, 0)
template <>
_CCCL_GLOBAL_CONSTANT __half e_v<__half> = __numbers<__half>::__e();
template <>
_CCCL_GLOBAL_CONSTANT __half log2e_v<__half> = __numbers<__half>::__log2e();
template <>
_CCCL_GLOBAL_CONSTANT __half log10e_v<__half> = __numbers<__half>::__log10e();
template <>
_CCCL_GLOBAL_CONSTANT __half pi_v<__half> = __numbers<__half>::__pi();
template <>
_CCCL_GLOBAL_CONSTANT __half inv_pi_v<__half> = __numbers<__half>::__inv_pi();
template <>
_CCCL_GLOBAL_CONSTANT __half inv_sqrtpi_v<__half> = __numbers<__half>::__inv_sqrtpi();
template <>
_CCCL_GLOBAL_CONSTANT __half ln2_v<__half> = __numbers<__half>::__ln2();
template <>
_CCCL_GLOBAL_CONSTANT __half ln10_v<__half> = __numbers<__half>::__ln10();
template <>
_CCCL_GLOBAL_CONSTANT __half sqrt2_v<__half> = __numbers<__half>::__sqrt2();
template <>
_CCCL_GLOBAL_CONSTANT __half sqrt3_v<__half> = __numbers<__half>::__sqrt3();
template <>
_CCCL_GLOBAL_CONSTANT __half inv_sqrt3_v<__half> = __numbers<__half>::__inv_sqrt3();
template <>
_CCCL_GLOBAL_CONSTANT __half egamma_v<__half> = __numbers<__half>::__egamma();
template <>
_CCCL_GLOBAL_CONSTANT __half phi_v<__half> = __numbers<__half>::__phi();
# endif // LIBHIPCXX_ROCM_VERSION_GE(7, 11, 0)
# endif // _LIBCUDACXX_HAS_NVFP16()
# if _LIBCUDACXX_HAS_NVBF16()
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 e_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__e();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 log2e_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__log2e();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 log10e_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__log10e();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 pi_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__pi();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 inv_pi_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__inv_pi();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 inv_sqrtpi_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__inv_sqrtpi();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 ln2_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__ln2();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 ln10_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__ln10();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 sqrt2_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__sqrt2();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 sqrt3_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__sqrt3();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 inv_sqrt3_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__inv_sqrt3();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 egamma_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__egamma();
template <>
_CCCL_GLOBAL_CONSTANT __hip_bfloat16 phi_v<__hip_bfloat16> = __numbers<__hip_bfloat16>::__phi();
# endif // _LIBCUDACXX_HAS_NVBF16()
#endif // !_CCCL_COMPILER(MSVC)
inline constexpr double e = __numbers<double>::__e();
inline constexpr double log2e = __numbers<double>::__log2e();
inline constexpr double log10e = __numbers<double>::__log10e();
inline constexpr double pi = __numbers<double>::__pi();
inline constexpr double inv_pi = __numbers<double>::__inv_pi();
inline constexpr double inv_sqrtpi = __numbers<double>::__inv_sqrtpi();
inline constexpr double ln2 = __numbers<double>::__ln2();
inline constexpr double ln10 = __numbers<double>::__ln10();
inline constexpr double sqrt2 = __numbers<double>::__sqrt2();
inline constexpr double sqrt3 = __numbers<double>::__sqrt3();
inline constexpr double inv_sqrt3 = __numbers<double>::__inv_sqrt3();
inline constexpr double egamma = __numbers<double>::__egamma();
inline constexpr double phi = __numbers<double>::__phi();
} // namespace numbers
_LIBCUDACXX_END_NAMESPACE_STD
#endif // _CUDA_STD_NUMBERS