blob: 2b2167e6ad331a7c11bc1ec8ae50ff99c6aa3a2a (
plain) (
blame)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
|
//===----------------------------------------------------------------------===//
//
// 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___CMATH_ISNORMAL_H
#define _LIBCUDACXX___CMATH_ISNORMAL_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/__cmath/fpclassify.h>
#include <cuda/std/__concepts/concept_macros.h>
#include <cuda/std/__floating_point/nvfp_types.h>
#include <cuda/std/__type_traits/is_integral.h>
#include <cuda/std/__cccl/prologue.h>
_LIBCUDACXX_BEGIN_NAMESPACE_STD
#if _CCCL_CHECK_BUILTIN(builtin_isnormal) || _CCCL_COMPILER(GCC)
# define _CCCL_BUILTIN_ISNORMAL(...) __builtin_isnormal(__VA_ARGS__)
#endif // _CCCL_CHECK_BUILTIN(isnormal)
// nvcc does not implement __builtin_isnormal
#if _CCCL_CUDA_COMPILER(NVCC)
# undef _CCCL_BUILTIN_ISNORMAL
#endif // _CCCL_CUDA_COMPILER(NVCC)
[[nodiscard]] _CCCL_API constexpr bool isnormal(float __x) noexcept
{
#if defined(_CCCL_BUILTIN_ISNORMAL)
return _CCCL_BUILTIN_ISNORMAL(__x);
#else // ^^^ _CCCL_BUILTIN_ISNORMAL ^^^ / vvv !_CCCL_BUILTIN_ISNORMAL vvv
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
#endif // !_CCCL_BUILTIN_ISNORMAL
}
[[nodiscard]] _CCCL_API constexpr bool isnormal(double __x) noexcept
{
#if defined(_CCCL_BUILTIN_ISNORMAL)
return _CCCL_BUILTIN_ISNORMAL(__x);
#else // ^^^ _CCCL_BUILTIN_ISNORMAL ^^^ / vvv !_CCCL_BUILTIN_ISNORMAL vvv
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
#endif // !_CCCL_BUILTIN_ISNORMAL
}
#if _CCCL_HAS_LONG_DOUBLE()
[[nodiscard]] _CCCL_API constexpr bool isnormal(long double __x) noexcept
{
# if defined(_CCCL_BUILTIN_ISNORMAL)
return _CCCL_BUILTIN_ISNORMAL(__x);
# else // ^^^ _CCCL_BUILTIN_ISNORMAL ^^^ / vvv !_CCCL_BUILTIN_ISNORMAL vvv
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
# endif // !_CCCL_BUILTIN_ISNORMAL
}
#endif // _CCCL_HAS_LONG_DOUBLE()
#if _CCCL_HAS_NVFP16()
[[nodiscard]] _CCCL_API constexpr bool isnormal(__half __x) noexcept
{
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
}
#endif // _CCCL_HAS_NVFP16()
#if _CCCL_HAS_NVBF16()
[[nodiscard]] _CCCL_API constexpr bool isnormal(__nv_bfloat16 __x) noexcept
{
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
}
#endif // _CCCL_HAS_NVBF16()
#if _CCCL_HAS_NVFP8_E4M3()
[[nodiscard]] _CCCL_API constexpr bool isnormal(__nv_fp8_e4m3 __x) noexcept
{
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
}
#endif // _CCCL_HAS_NVFP8_E4M3()
#if _CCCL_HAS_NVFP8_E5M2()
[[nodiscard]] _CCCL_API constexpr bool isnormal(__nv_fp8_e5m2 __x) noexcept
{
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
}
#endif // _CCCL_HAS_NVFP8_E5M2()
#if _CCCL_HAS_NVFP8_E8M0()
[[nodiscard]] _CCCL_API constexpr bool isnormal(__nv_fp8_e8m0 __x) noexcept
{
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
}
#endif // _CCCL_HAS_NVFP8_E8M0()
#if _CCCL_HAS_NVFP6_E2M3()
[[nodiscard]] _CCCL_API constexpr bool isnormal(__nv_fp6_e2m3 __x) noexcept
{
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
}
#endif // _CCCL_HAS_NVFP6_E2M3()
#if _CCCL_HAS_NVFP6_E3M2()
[[nodiscard]] _CCCL_API constexpr bool isnormal(__nv_fp6_e3m2 __x) noexcept
{
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
}
#endif // _CCCL_HAS_NVFP6_E3M2()
#if _CCCL_HAS_NVFP4_E2M1()
[[nodiscard]] _CCCL_API constexpr bool isnormal(__nv_fp4_e2m1 __x) noexcept
{
return _CUDA_VSTD::fpclassify(__x) == FP_NORMAL;
}
#endif // _CCCL_HAS_NVFP4_E2M1()
_CCCL_TEMPLATE(class _Tp)
_CCCL_REQUIRES(_CCCL_TRAIT(is_integral, _Tp))
[[nodiscard]] _CCCL_API constexpr bool isnormal(_Tp __x) noexcept
{
return __x != 0;
}
_LIBCUDACXX_END_NAMESPACE_STD
#include <cuda/std/__cccl/epilogue.h>
#endif // _LIBCUDACXX___CMATH_ISNORMAL_H
|
