Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
102 changes: 102 additions & 0 deletions libcudacxx/include/cuda/std/__floating_point/constants.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,102 @@
//===----------------------------------------------------------------------===//
//
// 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 _LIBCUDACXX___FLOATING_POINT_CONSTANTS_H
#define _LIBCUDACXX___FLOATING_POINT_CONSTANTS_H

#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/__concepts/concept_macros.h>
#include <cuda/std/__floating_point/format.h>
#include <cuda/std/__floating_point/mask.h>
#include <cuda/std/__floating_point/properties.h>
#include <cuda/std/__floating_point/storage.h>

_LIBCUDACXX_BEGIN_NAMESPACE_STD

template <__fp_format _Fmt>
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_t<_Fmt> __fp_inf() noexcept
{
static_assert(__fp_has_inf_v<_Fmt>, "The format does not support infinity");

return __fp_exp_mask_v<_Fmt>;
}

template <class _Tp>
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_inf() noexcept
{
return _CUDA_VSTD::__fp_from_storage<_Tp>(_CUDA_VSTD::__fp_inf<__fp_format_of_v<_Tp>>());
}

template <__fp_format _Fmt>
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_t<_Fmt> __fp_nan() noexcept
{
static_assert(__fp_has_nan_v<_Fmt>, "The format does not support nan");

if constexpr (_Fmt == __fp_format::__fp8_nv_e4m3)
{
return __fp_storage_t<_Fmt>(0x7fu);
}
else if constexpr (_Fmt == __fp_format::__fp8_nv_e8m0)
{
return __fp_storage_t<_Fmt>(0xffu);
}
else if constexpr (__fp_has_implicit_bit_v<_Fmt>)
{
return static_cast<__fp_storage_t<_Fmt>>(
__fp_exp_mask_v<_Fmt> | (__fp_storage_t<_Fmt>(1) << (__fp_mant_nbits_v<_Fmt> - 1)));
}
else
{
return static_cast<__fp_storage_t<_Fmt>>(
__fp_exp_mask_v<_Fmt> | (__fp_storage_t<_Fmt>(3) << (__fp_mant_nbits_v<_Fmt> - 2)));
}
}

template <class _Tp>
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_nan() noexcept
{
return _CUDA_VSTD::__fp_from_storage<_Tp>(_CUDA_VSTD::__fp_nan<__fp_format_of_v<_Tp>>());
}

template <__fp_format _Fmt>
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_t<_Fmt> __fp_nans() noexcept
{
static_assert(__fp_has_nans_v<_Fmt>, "The format does not support nans");

if constexpr (__fp_has_implicit_bit_v<_Fmt>)
{
return static_cast<__fp_storage_t<_Fmt>>(
__fp_exp_mask_v<_Fmt> | (__fp_storage_t<_Fmt>(1) << (__fp_mant_nbits_v<_Fmt> - 2)));
}
else
{
return static_cast<__fp_storage_t<_Fmt>>(
__fp_exp_mask_v<_Fmt> | (__fp_storage_t<_Fmt>(5) << (__fp_mant_nbits_v<_Fmt> - 3)));
}
}

template <class _Tp>
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_nans() noexcept
{
return _CUDA_VSTD::__fp_from_storage<_Tp>(_CUDA_VSTD::__fp_nans<__fp_format_of_v<_Tp>>());
}

_LIBCUDACXX_END_NAMESPACE_STD

#endif // _LIBCUDACXX___FLOATING_POINT_CONSTANTS_H
1 change: 1 addition & 0 deletions libcudacxx/include/cuda/std/__floating_point/fp.h
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,7 @@
#include <cuda/std/__floating_point/cast.h>
#include <cuda/std/__floating_point/cccl_fp.h>
#include <cuda/std/__floating_point/common_type.h>
#include <cuda/std/__floating_point/constants.h>
#include <cuda/std/__floating_point/conversion_rank_order.h>
#include <cuda/std/__floating_point/format.h>
#include <cuda/std/__floating_point/mask.h>
Expand Down
54 changes: 54 additions & 0 deletions libcudacxx/include/cuda/std/__floating_point/properties.h
Original file line number Diff line number Diff line change
Expand Up @@ -166,6 +166,60 @@ inline constexpr bool __fp_has_denorm_v = true;
template <>
inline constexpr bool __fp_has_denorm_v<__fp_format::__fp8_nv_e8m0> = false;

// __fp_has_inf_v

template <__fp_format _Fmt>
inline constexpr bool __fp_has_inf_v = true;

template <>
inline constexpr bool __fp_has_inf_v<__fp_format::__fp8_nv_e4m3> = false;

template <>
inline constexpr bool __fp_has_inf_v<__fp_format::__fp8_nv_e8m0> = false;

template <>
inline constexpr bool __fp_has_inf_v<__fp_format::__fp6_nv_e2m3> = false;

template <>
inline constexpr bool __fp_has_inf_v<__fp_format::__fp6_nv_e3m2> = false;

template <>
inline constexpr bool __fp_has_inf_v<__fp_format::__fp4_nv_e2m1> = false;

// __fp_has_nan_v

template <__fp_format _Fmt>
inline constexpr bool __fp_has_nan_v = true;

template <>
inline constexpr bool __fp_has_nan_v<__fp_format::__fp6_nv_e2m3> = false;

template <>
inline constexpr bool __fp_has_nan_v<__fp_format::__fp6_nv_e3m2> = false;

template <>
inline constexpr bool __fp_has_nan_v<__fp_format::__fp4_nv_e2m1> = false;

// __fp_has_nans_v

template <__fp_format _Fmt>
inline constexpr bool __fp_has_nans_v = true;

template <>
inline constexpr bool __fp_has_nans_v<__fp_format::__fp8_nv_e4m3> = false;

template <>
inline constexpr bool __fp_has_nans_v<__fp_format::__fp8_nv_e8m0> = false;

template <>
inline constexpr bool __fp_has_nans_v<__fp_format::__fp6_nv_e2m3> = false;

template <>
inline constexpr bool __fp_has_nans_v<__fp_format::__fp6_nv_e3m2> = false;

template <>
inline constexpr bool __fp_has_nans_v<__fp_format::__fp4_nv_e2m1> = false;

_LIBCUDACXX_END_NAMESPACE_STD

#endif // _LIBCUDACXX___FLOATING_POINT_PROPERTIES_H
Original file line number Diff line number Diff line change
@@ -0,0 +1,88 @@
//===----------------------------------------------------------------------===//
//
// 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.
//
//===----------------------------------------------------------------------===//

#include <cuda/std/__floating_point/fp.h>
#include <cuda/std/cassert>
#include <cuda/std/cmath>
#include <cuda/std/cstring>
#include <cuda/std/limits>
#include <cuda/std/type_traits>

#include "test_macros.h"

template <class T>
__host__ __device__ void test_fp_storage()
{
constexpr auto fmt = cuda::std::__fp_format_of_v<T>;

// __fp_has_inf_v must match numeric_limits::has_infinity
static_assert(cuda::std::__fp_has_inf_v<fmt> == cuda::std::numeric_limits<T>::has_infinity);

// test __fp_inf value to match numeric_limits::infinity()
if constexpr (cuda::std::__fp_has_inf_v<fmt>)
{
const auto val = cuda::std::__fp_inf<T>();
const auto ref = cuda::std::numeric_limits<T>::infinity();
assert(cuda::std::memcmp(&val, &ref, sizeof(T)) == 0);
}

// __fp_has_nan_v must match numeric_limits::has_quiet_NaN
static_assert(cuda::std::__fp_has_nan_v<fmt> == cuda::std::numeric_limits<T>::has_quiet_NaN);

// test __fp_nan value
if constexpr (cuda::std::__fp_has_nan_v<fmt>)
{
assert(cuda::std::isnan(cuda::std::__fp_nan<T>()));
}

// __fp_has_nans_v must match numeric_limits::has_signaling_NaN
static_assert(cuda::std::__fp_has_nans_v<fmt> == cuda::std::numeric_limits<T>::has_signaling_NaN);

// test __fp_nans value
if constexpr (cuda::std::__fp_has_nans_v<fmt>)
{
assert(cuda::std::isnan(cuda::std::__fp_nans<T>()));
}
}

int main(int, char**)
{
test_fp_storage<float>();
test_fp_storage<double>();
#if _CCCL_HAS_LONG_DOUBLE()
test_fp_storage<long double>();
#endif // _CCCL_HAS_LONG_DOUBLE()
#if _CCCL_HAS_NVFP16()
test_fp_storage<__half>();
#endif // _CCCL_HAS_NVFP16()
#if _CCCL_HAS_NVBF16()
test_fp_storage<__nv_bfloat16>();
#endif // _CCCL_HAS_NVBF16()
#if _CCCL_HAS_NVFP8_E4M3()
test_fp_storage<__nv_fp8_e4m3>();
#endif // _CCCL_HAS_NVFP8_E4M3()
#if _CCCL_HAS_NVFP8_E5M2()
test_fp_storage<__nv_fp8_e5m2>();
#endif // _CCCL_HAS_NVFP8_E5M2()
#if _CCCL_HAS_NVFP8_E8M0()
test_fp_storage<__nv_fp8_e8m0>();
#endif // _CCCL_HAS_NVFP8_E8M0()
#if _CCCL_HAS_NVFP6_E2M3()
test_fp_storage<__nv_fp6_e2m3>();
#endif // _CCCL_HAS_NVFP6_E2M3()
#if _CCCL_HAS_NVFP6_E3M2()
test_fp_storage<__nv_fp6_e3m2>();
#endif // _CCCL_HAS_NVFP6_E3M2()
#if _CCCL_HAS_NVFP4_E2M1()
test_fp_storage<__nv_fp4_e2m1>();
#endif // _CCCL_HAS_NVFP4_E2M1()

return 0;
}
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__binary16> =
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__binary16>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__binary16> == 11);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__binary16>);
static_assert(cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__binary16>);
static_assert(cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__binary16>);
static_assert(cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__binary16>);

// binary32

Expand All @@ -40,6 +43,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__binary32> =
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__binary32>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__binary32> == 24);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__binary32>);
static_assert(cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__binary32>);
static_assert(cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__binary32>);
static_assert(cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__binary32>);

// binary64

Expand All @@ -52,6 +58,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__binary64> =
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__binary64>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__binary64> == 53);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__binary64>);
static_assert(cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__binary64>);
static_assert(cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__binary64>);
static_assert(cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__binary64>);

// binary128

Expand All @@ -64,6 +73,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__binary128>
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__binary128>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__binary128> == 113);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__binary128>);
static_assert(cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__binary128>);
static_assert(cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__binary128>);
static_assert(cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__binary128>);

// bfloat16

Expand All @@ -76,6 +88,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__bfloat16> =
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__bfloat16>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__bfloat16> == 8);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__bfloat16>);
static_assert(cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__bfloat16>);
static_assert(cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__bfloat16>);
static_assert(cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__bfloat16>);

// fp80_x86

Expand All @@ -88,6 +103,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__fp80_x86> =
static_assert(!cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__fp80_x86>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__fp80_x86> == 64);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__fp80_x86>);
static_assert(cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__fp80_x86>);
static_assert(cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__fp80_x86>);
static_assert(cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__fp80_x86>);

// fp8_nv_e4m3

Expand All @@ -100,6 +118,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__fp8_nv_e4m3
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__fp8_nv_e4m3>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__fp8_nv_e4m3> == 4);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__fp8_nv_e4m3>);
static_assert(!cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__fp8_nv_e4m3>);
static_assert(cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__fp8_nv_e4m3>);
static_assert(!cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__fp8_nv_e4m3>);

// fp8_nv_e5m2

Expand All @@ -112,6 +133,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__fp8_nv_e5m2
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__fp8_nv_e5m2>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__fp8_nv_e5m2> == 3);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__fp8_nv_e5m2>);
static_assert(cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__fp8_nv_e5m2>);
static_assert(cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__fp8_nv_e5m2>);
static_assert(cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__fp8_nv_e5m2>);

// fp8_nv_e8m0

Expand All @@ -124,6 +148,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__fp8_nv_e8m0
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__fp8_nv_e8m0>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__fp8_nv_e8m0> == 1);
static_assert(!cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__fp8_nv_e8m0>);
static_assert(!cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__fp8_nv_e8m0>);
static_assert(cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__fp8_nv_e8m0>);
static_assert(!cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__fp8_nv_e8m0>);

// fp6_nv_e2m3

Expand All @@ -136,6 +163,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__fp6_nv_e2m3
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__fp6_nv_e2m3>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__fp6_nv_e2m3> == 4);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__fp6_nv_e2m3>);
static_assert(!cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__fp6_nv_e2m3>);
static_assert(!cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__fp6_nv_e2m3>);
static_assert(!cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__fp6_nv_e2m3>);

// fp6_nv_e3m2

Expand All @@ -148,6 +178,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__fp6_nv_e3m2
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__fp6_nv_e3m2>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__fp6_nv_e3m2> == 3);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__fp6_nv_e3m2>);
static_assert(!cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__fp6_nv_e3m2>);
static_assert(!cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__fp6_nv_e3m2>);
static_assert(!cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__fp6_nv_e3m2>);

// fp4_nv_e2m1

Expand All @@ -160,6 +193,9 @@ static_assert(cuda::std::__fp_mant_nbits_v<cuda::std::__fp_format::__fp4_nv_e2m1
static_assert(cuda::std::__fp_has_implicit_bit_v<cuda::std::__fp_format::__fp4_nv_e2m1>);
static_assert(cuda::std::__fp_digits_v<cuda::std::__fp_format::__fp4_nv_e2m1> == 2);
static_assert(cuda::std::__fp_has_denorm_v<cuda::std::__fp_format::__fp4_nv_e2m1>);
static_assert(!cuda::std::__fp_has_inf_v<cuda::std::__fp_format::__fp4_nv_e2m1>);
static_assert(!cuda::std::__fp_has_nan_v<cuda::std::__fp_format::__fp4_nv_e2m1>);
static_assert(!cuda::std::__fp_has_nans_v<cuda::std::__fp_format::__fp4_nv_e2m1>);

int main(int, char**)
{
Expand Down
Loading