contrib/libs/nvidia/cccl/libcudacxx/include/cuda/std/__cmath/hyperbolic_functions.h


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
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285

// -*- C++ -*-
//===----------------------------------------------------------------------===//
//
// 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.
//
//===----------------------------------------------------------------------===//

#ifndef _LIBCUDACXX___CMATH_HYPERBOLIC_FUNCTIONS_H
#define _LIBCUDACXX___CMATH_HYPERBOLIC_FUNCTIONS_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/__floating_point/nvfp_types.h>
#include <cuda/std/__type_traits/enable_if.h>
#include <cuda/std/__type_traits/is_integral.h>

#include <nv/target>

#if !_CCCL_COMPILER(NVRTC)
#  include <math.h>
#endif // !_CCCL_COMPILER(NVRTC)

#include <cuda/std/__cccl/prologue.h>

_LIBCUDACXX_BEGIN_NAMESPACE_STD

// cosh

#if _CCCL_CHECK_BUILTIN(builtin_cosh) || _CCCL_COMPILER(GCC)
#  define _CCCL_BUILTIN_COSHF(...) __builtin_coshf(__VA_ARGS__)
#  define _CCCL_BUILTIN_COSH(...)  __builtin_cosh(__VA_ARGS__)
#  define _CCCL_BUILTIN_COSHL(...) __builtin_coshl(__VA_ARGS__)
#endif // _CCCL_CHECK_BUILTIN(builtin_cosh)

#if _CCCL_CUDA_COMPILER(CLANG)
#  undef _CCCL_BUILTIN_COSHF
#  undef _CCCL_BUILTIN_COSH
#  undef _CCCL_BUILTIN_COSHL
#endif // _CCCL_CUDA_COMPILER(CLANG)

[[nodiscard]] _CCCL_API inline float cosh(float __x) noexcept
{
#if defined(_CCCL_BUILTIN_COSHF)
  return _CCCL_BUILTIN_COSHF(__x);
#else // ^^^ _CCCL_BUILTIN_COSHF ^^^ / vvv !_CCCL_BUILTIN_COSHF vvv
  return ::coshf(__x);
#endif // !_CCCL_BUILTIN_COSHF
}

[[nodiscard]] _CCCL_API inline float coshf(float __x) noexcept
{
#if defined(_CCCL_BUILTIN_COSHF)
  return _CCCL_BUILTIN_COSHF(__x);
#else // ^^^ _CCCL_BUILTIN_COSHF ^^^ / vvv !_CCCL_BUILTIN_COSHF vvv
  return ::coshf(__x);
#endif // !_CCCL_BUILTIN_COSHF
}

[[nodiscard]] _CCCL_API inline double cosh(double __x) noexcept
{
#if defined(_CCCL_BUILTIN_COSH)
  return _CCCL_BUILTIN_COSH(__x);
#else // ^^^ _CCCL_BUILTIN_COSH ^^^ / vvv !_CCCL_BUILTIN_COSH vvv
  return ::cosh(__x);
#endif // !_CCCL_BUILTIN_COSH
}

#if _CCCL_HAS_LONG_DOUBLE()
[[nodiscard]] _CCCL_API inline long double cosh(long double __x) noexcept
{
#  if defined(_CCCL_BUILTIN_COSHL)
  return _CCCL_BUILTIN_COSHL(__x);
#  else // ^^^ _CCCL_BUILTIN_COSHL ^^^ / vvv !_CCCL_BUILTIN_COSHL vvv
  return ::coshl(__x);
#  endif // !_CCCL_BUILTIN_COSHL
}

[[nodiscard]] _CCCL_API inline long double coshl(long double __x) noexcept
{
#  if defined(_CCCL_BUILTIN_COSHL)
  return _CCCL_BUILTIN_COSHL(__x);
#  else // ^^^ _CCCL_BUILTIN_COSHL ^^^ / vvv !_CCCL_BUILTIN_COSHL vvv
  return ::coshl(__x);
#  endif // !_CCCL_BUILTIN_COSHL
}
#endif // _CCCL_HAS_LONG_DOUBLE()

#if _LIBCUDACXX_HAS_NVFP16()
[[nodiscard]] _CCCL_API inline __half cosh(__half __x) noexcept
{
  return __float2half(_CUDA_VSTD::coshf(__half2float(__x)));
}
#endif // _LIBCUDACXX_HAS_NVFP16()

#if _LIBCUDACXX_HAS_NVBF16()
[[nodiscard]] _CCCL_API inline __nv_bfloat16 cosh(__nv_bfloat16 __x) noexcept
{
  return __float2bfloat16(_CUDA_VSTD::coshf(__bfloat162float(__x)));
}
#endif // _LIBCUDACXX_HAS_NVBF16()

template <class _Integer, enable_if_t<_CCCL_TRAIT(is_integral, _Integer), int> = 0>
[[nodiscard]] _CCCL_API inline double cosh(_Integer __x) noexcept
{
  return _CUDA_VSTD::cosh((double) __x);
}

// sinh

#if _CCCL_CHECK_BUILTIN(builtin_sinh) || _CCCL_COMPILER(GCC)
#  define _CCCL_BUILTIN_SINHF(...) __builtin_sinhf(__VA_ARGS__)
#  define _CCCL_BUILTIN_SINH(...)  __builtin_sinh(__VA_ARGS__)
#  define _CCCL_BUILTIN_SINHL(...) __builtin_sinhl(__VA_ARGS__)
#endif // _CCCL_CHECK_BUILTIN(builtin_sin)

#if _CCCL_CUDA_COMPILER(CLANG)
#  undef _CCCL_BUILTIN_SINHF
#  undef _CCCL_BUILTIN_SINH
#  undef _CCCL_BUILTIN_SINHL
#endif // _CCCL_CUDA_COMPILER(CLANG)

[[nodiscard]] _CCCL_API inline float sinh(float __x) noexcept
{
#if defined(_CCCL_BUILTIN_SINHF)
  return _CCCL_BUILTIN_SINHF(__x);
#else // ^^^ _CCCL_BUILTIN_SINHF ^^^ / vvv !_CCCL_BUILTIN_SINHF vvv
  return ::sinhf(__x);
#endif // !_CCCL_BUILTIN_SINHF
}

[[nodiscard]] _CCCL_API inline float sinhf(float __x) noexcept
{
#if defined(_CCCL_BUILTIN_SINHF)
  return _CCCL_BUILTIN_SINHF(__x);
#else // ^^^ _CCCL_BUILTIN_SINHF ^^^ / vvv !_CCCL_BUILTIN_SINHF vvv
  return ::sinhf(__x);
#endif // !_CCCL_BUILTIN_SINHF
}

[[nodiscard]] _CCCL_API inline double sinh(double __x) noexcept
{
#if defined(_CCCL_BUILTIN_SINH)
  return _CCCL_BUILTIN_SINH(__x);
#else // ^^^ _CCCL_BUILTIN_SINH ^^^ / vvv !_CCCL_BUILTIN_SINH vvv
  return ::sinh(__x);
#endif // !_CCCL_BUILTIN_SINH
}

#if _CCCL_HAS_LONG_DOUBLE()
[[nodiscard]] _CCCL_API inline long double sinh(long double __x) noexcept
{
#  if defined(_CCCL_BUILTIN_SINHL)
  return _CCCL_BUILTIN_SINHL(__x);
#  else // ^^^ _CCCL_BUILTIN_SINHL ^^^ / vvv !_CCCL_BUILTIN_SINHL vvv
  return ::sinhl(__x);
#  endif // !_CCCL_BUILTIN_SINHL
}

[[nodiscard]] _CCCL_API inline long double sinhl(long double __x) noexcept
{
#  if defined(_CCCL_BUILTIN_SINHL)
  return _CCCL_BUILTIN_SINHL(__x);
#  else // ^^^ _CCCL_BUILTIN_SINHL ^^^ / vvv !_CCCL_BUILTIN_SINHL vvv
  return ::sinhl(__x);
#  endif // !_CCCL_BUILTIN_SINHL
}
#endif // _CCCL_HAS_LONG_DOUBLE()

#if _LIBCUDACXX_HAS_NVFP16()
[[nodiscard]] _CCCL_API inline __half sinh(__half __x) noexcept
{
  return __float2half(_CUDA_VSTD::sinhf(__half2float(__x)));
}
#endif // _LIBCUDACXX_HAS_NVFP16()

#if _LIBCUDACXX_HAS_NVBF16()
[[nodiscard]] _CCCL_API inline __nv_bfloat16 sinh(__nv_bfloat16 __x) noexcept
{
  return __float2bfloat16(_CUDA_VSTD::sinhf(__bfloat162float(__x)));
}
#endif // _LIBCUDACXX_HAS_NVBF16()

template <class _Integer, enable_if_t<_CCCL_TRAIT(is_integral, _Integer), int> = 0>
[[nodiscard]] _CCCL_API inline double sinh(_Integer __x) noexcept
{
  return _CUDA_VSTD::sinh((double) __x);
}

// tanh

#if _CCCL_CHECK_BUILTIN(builtin_tanh) || _CCCL_COMPILER(GCC)
#  define _CCCL_BUILTIN_TANHF(...) __builtin_tanhf(__VA_ARGS__)
#  define _CCCL_BUILTIN_TANH(...)  __builtin_tanh(__VA_ARGS__)
#  define _CCCL_BUILTIN_TANHL(...) __builtin_tanhl(__VA_ARGS__)
#endif // _CCCL_CHECK_BUILTIN(builtin_tan)

#if _CCCL_CUDA_COMPILER(CLANG)
#  undef _CCCL_BUILTIN_TANHF
#  undef _CCCL_BUILTIN_TANH
#  undef _CCCL_BUILTIN_TANHL
#endif // _CCCL_CUDA_COMPILER(CLANG)

[[nodiscard]] _CCCL_API inline float tanh(float __x) noexcept
{
#if defined(_CCCL_BUILTIN_TANHF)
  return _CCCL_BUILTIN_TANHF(__x);
#else // ^^^ _CCCL_BUILTIN_TANHF ^^^ / vvv !_CCCL_BUILTIN_TANHF vvv
  return ::tanhf(__x);
#endif // !_CCCL_BUILTIN_TANHF
}

[[nodiscard]] _CCCL_API inline float tanhf(float __x) noexcept
{
#if defined(_CCCL_BUILTIN_TANHF)
  return _CCCL_BUILTIN_TANHF(__x);
#else // ^^^ _CCCL_BUILTIN_TANHF ^^^ / vvv !_CCCL_BUILTIN_TANHF vvv
  return ::tanhf(__x);
#endif // !_CCCL_BUILTIN_TANHF
}

[[nodiscard]] _CCCL_API inline double tanh(double __x) noexcept
{
#if defined(_CCCL_BUILTIN_TANH)
  return _CCCL_BUILTIN_TANH(__x);
#else // ^^^ _CCCL_BUILTIN_TANH ^^^ / vvv !_CCCL_BUILTIN_TANH vvv
  return ::tanh(__x);
#endif // !_CCCL_BUILTIN_TANH
}

#if _CCCL_HAS_LONG_DOUBLE()
[[nodiscard]] _CCCL_API inline long double tanh(long double __x) noexcept
{
#  if defined(_CCCL_BUILTIN_TANHL)
  return _CCCL_BUILTIN_TANHL(__x);
#  else // ^^^ _CCCL_BUILTIN_TANHL ^^^ / vvv !_CCCL_BUILTIN_TANHL vvv
  return ::tanhl(__x);
#  endif // !_CCCL_BUILTIN_TANHL
}

[[nodiscard]] _CCCL_API inline long double tanhl(long double __x) noexcept
{
#  if defined(_CCCL_BUILTIN_TANHL)
  return _CCCL_BUILTIN_TANHL(__x);
#  else // ^^^ _CCCL_BUILTIN_TANHL ^^^ / vvv !_CCCL_BUILTIN_TANHL vvv
  return ::tanhl(__x);
#  endif // !_CCCL_BUILTIN_TANHL
}
#endif // _CCCL_HAS_LONG_DOUBLE()

#if _LIBCUDACXX_HAS_NVFP16()
[[nodiscard]] _CCCL_API inline __half tanh(__half __x) noexcept
{
  return __float2half(_CUDA_VSTD::tanhf(__half2float(__x)));
}
#endif // _LIBCUDACXX_HAS_NVFP16()

#if _LIBCUDACXX_HAS_NVBF16()
[[nodiscard]] _CCCL_API inline __nv_bfloat16 tanh(__nv_bfloat16 __x) noexcept
{
  return __float2bfloat16(_CUDA_VSTD::tanhf(__bfloat162float(__x)));
}
#endif // _LIBCUDACXX_HAS_NVBF16()

template <class _Integer, enable_if_t<_CCCL_TRAIT(is_integral, _Integer), int> = 0>
[[nodiscard]] _CCCL_API inline double tanh(_Integer __x) noexcept
{
  return _CUDA_VSTD::tanh((double) __x);
}

_LIBCUDACXX_END_NAMESPACE_STD

#include <cuda/std/__cccl/epilogue.h>

#endif // _LIBCUDACXX___CMATH_HYPERBOLIC_FUNCTIONS_H