codekingpro/portable-devtools
114k
1/*
2* Copyright 1993-2024 NVIDIA Corporation. All rights reserved.
3*
4* NOTICE TO LICENSEE:
5*
6* This source code and/or documentation ("Licensed Deliverables") are
7* subject to NVIDIA intellectual property rights under U.S. and
8* international Copyright laws.
9*
10* These Licensed Deliverables contained herein is PROPRIETARY and
11* CONFIDENTIAL to NVIDIA and is being provided under the terms and
12* conditions of a form of NVIDIA software license agreement by and
13* between NVIDIA and Licensee ("License Agreement") or electronically
14* accepted by Licensee. Notwithstanding any terms or conditions to
15* the contrary in the License Agreement, reproduction or disclosure
16* of the Licensed Deliverables to any third party without the express
17* written consent of NVIDIA is prohibited.
18*
19* NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
20* LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE
21* SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS
22* PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND.
23* NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED
24* DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY,
25* NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
26* NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
27* LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY
28* SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY
29* DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,
30* WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS
31* ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE
32* OF THESE LICENSED DELIVERABLES.
33*
34* U.S. Government End Users. These Licensed Deliverables are a
35* "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT
36* 1995), consisting of "commercial computer software" and "commercial
37* computer software documentation" as such terms are used in 48
38* C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government
39* only as a commercial end item. Consistent with 48 C.F.R.12.212 and
40* 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all
41* U.S. Government End Users acquire the Licensed Deliverables with
42* only those rights set forth herein.
43*
44* Any use of the Licensed Deliverables in individual and commercial
45* software must include, in the user documentation and internal
46* comments to the code, the above Disclaimer and U.S. Government End
47* Users Notice.
48*/
49
50/**
51* \defgroup CUDA_MATH_INTRINSIC_BFLOAT16 Bfloat16 Precision Intrinsics
52* This section describes nv_bfloat16 precision intrinsic functions.
53* To use these functions, include the header file \p cuda_bf16.h in your program.
54* All of the functions defined here are available in device code.
55* Some of the functions are also available to host compilers, please
56* refer to respective functions' documentation for details.
57*
58* NOTE: Aggressive floating-point optimizations performed by host or device
59* compilers may affect numeric behavior of the functions implemented in this
60* header. Specific examples are:
61* - hsin(__nv_bfloat16);
62* - hcos(__nv_bfloat16);
63* - h2sin(__nv_bfloat162);
64* - h2cos(__nv_bfloat162);
65*
66* The following macros are available to help users selectively enable/disable
67* various definitions present in the header file:
68* - \p CUDA_NO_BFLOAT16 - If defined, this macro will prevent the definition of
69* additional type aliases in the global namespace, helping to avoid potential
70* conflicts with symbols defined in the user program.
71* - \p __CUDA_NO_BFLOAT16_CONVERSIONS__ - If defined, this macro will prevent
72* the use of the C++ type conversions (converting constructors and conversion
73* operators) that are common for built-in floating-point types, but may be
74* undesirable for \p __nv_bfloat16 which is essentially a user-defined type.
75* - \p __CUDA_NO_BFLOAT16_OPERATORS__ and \p __CUDA_NO_BFLOAT162_OPERATORS__ -
76* If defined, these macros will prevent the inadvertent use of usual arithmetic
77* and comparison operators. This enforces the storage-only type semantics and
78* prevents C++ style computations on \p __nv_bfloat16 and \p __nv_bfloat162 types.
79*/
80
81/**
82* \defgroup CUDA_MATH_INTRINSIC_BFLOAT16_CONSTANTS Bfloat16 Arithmetic Constants
83* \ingroup CUDA_MATH_INTRINSIC_BFLOAT16
84* To use these constants, include the header file \p cuda_bf16.h in your program.
85*/
86
87/**
88* \defgroup CUDA_MATH__BFLOAT16_ARITHMETIC Bfloat16 Arithmetic Functions
89* \ingroup CUDA_MATH_INTRINSIC_BFLOAT16
90* To use these functions, include the header file \p cuda_bf16.h in your program.
91*/
92
93/**
94* \defgroup CUDA_MATH__BFLOAT162_ARITHMETIC Bfloat162 Arithmetic Functions
95* \ingroup CUDA_MATH_INTRINSIC_BFLOAT16
96* To use these functions, include the header file \p cuda_bf16.h in your program.
97*/
98
99/**
100* \defgroup CUDA_MATH__BFLOAT16_COMPARISON Bfloat16 Comparison Functions
101* \ingroup CUDA_MATH_INTRINSIC_BFLOAT16
102* To use these functions, include the header file \p cuda_bf16.h in your program.
103*/
104
105/**
106* \defgroup CUDA_MATH__BFLOAT162_COMPARISON Bfloat162 Comparison Functions
107* \ingroup CUDA_MATH_INTRINSIC_BFLOAT16
108* To use these functions, include the header file \p cuda_bf16.h in your program.
109*/
110
111/**
112* \defgroup CUDA_MATH__BFLOAT16_MISC Bfloat16 Precision Conversion and Data Movement
113* \ingroup CUDA_MATH_INTRINSIC_BFLOAT16
114* To use these functions, include the header file \p cuda_bf16.h in your program.
115*/
116
117/**
118* \defgroup CUDA_MATH__BFLOAT16_FUNCTIONS Bfloat16 Math Functions
119* \ingroup CUDA_MATH_INTRINSIC_BFLOAT16
120* To use these functions, include the header file \p cuda_bf16.h in your program.
121*/
122
123/**
124* \defgroup CUDA_MATH__BFLOAT162_FUNCTIONS Bfloat162 Math Functions
125* \ingroup CUDA_MATH_INTRINSIC_BFLOAT16
126* To use these functions, include the header file \p cuda_bf16.h in your program.
127*/
128
129#ifndef __CUDA_BF16_H__
130#define __CUDA_BF16_H__
131
132/* bring in __half data type and operations, for use in converting constructors */
133#include "cuda_fp16.h"
134/* bring in float2, double4, etc vector types */
135#include "vector_types.h"
136/* bring in operations on vector types like: make_float2 */
137#include "vector_functions.h"
138
139#define ___CUDA_BF16_STRINGIFY_INNERMOST(x) #x
140#define __CUDA_BF16_STRINGIFY(x) ___CUDA_BF16_STRINGIFY_INNERMOST(x)
141
142#if defined(__cplusplus)
143
144/* Set up function decorations */
145#if (defined(__CUDACC_RTC__) && ((__CUDACC_VER_MAJOR__ > 12) || ((__CUDACC_VER_MAJOR__ == 12) && (__CUDACC_VER_MINOR__ >= 3))))
146#define __CUDA_BF16_DECL__ __device__
147#define __CUDA_HOSTDEVICE_BF16_DECL__ __device__
148#define __CUDA_HOSTDEVICE__ __device__
149#elif defined(__CUDACC__) || defined(_NVHPC_CUDA)
150#define __CUDA_BF16_DECL__ static __device__ __inline__
151#define __CUDA_HOSTDEVICE_BF16_DECL__ static __host__ __device__ __inline__
152#define __CUDA_HOSTDEVICE__ __host__ __device__
153#else /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
154#if defined(__GNUC__)
155#define __CUDA_HOSTDEVICE_BF16_DECL__ static __attribute__ ((unused))
156#else
157#define __CUDA_HOSTDEVICE_BF16_DECL__ static
158#endif /* defined(__GNUC__) */
159#define __CUDA_HOSTDEVICE__
160#endif /* (defined(__CUDACC_RTC__) && ((__CUDACC_VER_MAJOR__ > 12) || ((__CUDACC_VER_MAJOR__ == 12) && (__CUDACC_VER_MINOR__ >= 3))))) */
161
162#define __CUDA_BF16_TYPES_EXIST__
163
164/* Macros to allow nv_bfloat16 & nv_bfloat162 to be used by inline assembly */
165#define __BFLOAT16_TO_US(var) *(reinterpret_cast<unsigned short *>(&(var)))
166#define __BFLOAT16_TO_CUS(var) *(reinterpret_cast<const unsigned short *>(&(var)))
167#define __BFLOAT162_TO_UI(var) *(reinterpret_cast<unsigned int *>(&(var)))
168#define __BFLOAT162_TO_CUI(var) *(reinterpret_cast<const unsigned int *>(&(var)))
169
170/* Forward-declaration of structures defined in "cuda_bf16.hpp" */
171struct __nv_bfloat16;
172struct __nv_bfloat162;
173
174/**
175* \ingroup CUDA_MATH__BFLOAT16_MISC
176* \brief Converts double number to nv_bfloat16 precision in round-to-nearest-even mode
177* and returns \p nv_bfloat16 with converted value.
178*
179* \details Converts double number \p a to nv_bfloat16 precision in round-to-nearest-even mode.
180* \param[in] a - double. Is only being read.
181* \returns nv_bfloat16
182* - \p a converted to nv_bfloat16.
183* \internal
184* \exception-guarantee no-throw guarantee
185* \behavior reentrant, thread safe
186* \endinternal
187*/
188__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __double2bfloat16(const double a);
189/**
190* \ingroup CUDA_MATH__BFLOAT16_MISC
191* \brief Converts float number to nv_bfloat16 precision in round-to-nearest-even mode
192* and returns \p nv_bfloat16 with converted value.
193*
194* \details Converts float number \p a to nv_bfloat16 precision in round-to-nearest-even mode.
195* \param[in] a - float. Is only being read.
196* \returns nv_bfloat16
197* - \p a converted to nv_bfloat16.
198* \internal
199* \exception-guarantee no-throw guarantee
200* \behavior reentrant, thread safe
201* \endinternal
202*/
203__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __float2bfloat16(const float a);
204/**
205* \ingroup CUDA_MATH__BFLOAT16_MISC
206* \brief Converts float number to nv_bfloat16 precision in round-to-nearest-even mode
207* and returns \p nv_bfloat16 with converted value.
208*
209* \details Converts float number \p a to nv_bfloat16 precision in round-to-nearest-even mode.
210* \param[in] a - float. Is only being read.
211* \returns nv_bfloat16
212* - \p a converted to nv_bfloat16.
213* \internal
214* \exception-guarantee no-throw guarantee
215* \behavior reentrant, thread safe
216* \endinternal
217*/
218__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __float2bfloat16_rn(const float a);
219/**
220* \ingroup CUDA_MATH__BFLOAT16_MISC
221* \brief Converts float number to nv_bfloat16 precision in round-towards-zero mode
222* and returns \p nv_bfloat16 with converted value.
223*
224* \details Converts float number \p a to nv_bfloat16 precision in round-towards-zero mode.
225* \param[in] a - float. Is only being read.
226* \returns nv_bfloat16
227* - \p a converted to nv_bfloat16.
228* \internal
229* \exception-guarantee no-throw guarantee
230* \behavior reentrant, thread safe
231* \endinternal
232*/
233__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __float2bfloat16_rz(const float a);
234/**
235* \ingroup CUDA_MATH__BFLOAT16_MISC
236* \brief Converts float number to nv_bfloat16 precision in round-down mode
237* and returns \p nv_bfloat16 with converted value.
238*
239* \details Converts float number \p a to nv_bfloat16 precision in round-down mode.
240* \param[in] a - float. Is only being read.
241*
242* \returns nv_bfloat16
243* - \p a converted to nv_bfloat16.
244* \internal
245* \exception-guarantee no-throw guarantee
246* \behavior reentrant, thread safe
247* \endinternal
248*/
249__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __float2bfloat16_rd(const float a);
250/**
251* \ingroup CUDA_MATH__BFLOAT16_MISC
252* \brief Converts float number to nv_bfloat16 precision in round-up mode
253* and returns \p nv_bfloat16 with converted value.
254*
255* \details Converts float number \p a to nv_bfloat16 precision in round-up mode.
256* \param[in] a - float. Is only being read.
257*
258* \returns nv_bfloat16
259* - \p a converted to nv_bfloat16.
260* \internal
261* \exception-guarantee no-throw guarantee
262* \behavior reentrant, thread safe
263* \endinternal
264*/
265__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __float2bfloat16_ru(const float a);
266/**
267* \ingroup CUDA_MATH__BFLOAT16_MISC
268* \brief Converts \p nv_bfloat16 number to float.
269*
270* \details Converts nv_bfloat16 number \p a to float.
271* \param[in] a - float. Is only being read.
272*
273* \returns float
274* - \p a converted to float.
275* \internal
276* \exception-guarantee no-throw guarantee
277* \behavior reentrant, thread safe
278* \endinternal
279*/
280__CUDA_HOSTDEVICE_BF16_DECL__ float __bfloat162float(const __nv_bfloat16 a);
281/**
282* \ingroup CUDA_MATH__BFLOAT16_MISC
283* \brief Converts input to nv_bfloat16 precision in round-to-nearest-even mode and
284* populates both halves of \p nv_bfloat162 with converted value.
285*
286* \details Converts input \p a to nv_bfloat16 precision in round-to-nearest-even mode and
287* populates both halves of \p nv_bfloat162 with converted value.
288* \param[in] a - float. Is only being read.
289*
290* \returns nv_bfloat162
291* - The \p nv_bfloat162 value with both halves equal to the converted nv_bfloat16
292* precision number.
293* \internal
294* \exception-guarantee no-throw guarantee
295* \behavior reentrant, thread safe
296* \endinternal
297*/
298__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat162 __float2bfloat162_rn(const float a);
299/**
300* \ingroup CUDA_MATH__BFLOAT16_MISC
301* \brief Converts both input floats to nv_bfloat16 precision in round-to-nearest-even
302* mode and returns \p nv_bfloat162 with converted values.
303*
304* \details Converts both input floats to nv_bfloat16 precision in round-to-nearest-even mode
305* and combines the results into one \p nv_bfloat162 number. Low 16 bits of the return
306* value correspond to the input \p a, high 16 bits correspond to the input \p
307* b.
308* \param[in] a - float. Is only being read.
309* \param[in] b - float. Is only being read.
310*
311* \returns nv_bfloat162
312* - The \p nv_bfloat162 value with corresponding halves equal to the
313* converted input floats.
314* \internal
315* \exception-guarantee no-throw guarantee
316* \behavior reentrant, thread safe
317* \endinternal
318*/
319__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat162 __floats2bfloat162_rn(const float a, const float b);
320/**
321* \ingroup CUDA_MATH__BFLOAT16_MISC
322* \brief Converts low 16 bits of \p nv_bfloat162 to float and returns the result
323*
324* \details Converts low 16 bits of \p nv_bfloat162 input \p a to 32-bit floating-point number
325* and returns the result.
326* \param[in] a - nv_bfloat162. Is only being read.
327*
328* \returns float
329* - The low 16 bits of \p a converted to float.
330* \internal
331* \exception-guarantee no-throw guarantee
332* \behavior reentrant, thread safe
333* \endinternal
334*/
335__CUDA_HOSTDEVICE_BF16_DECL__ float __low2float(const __nv_bfloat162 a);
336/**
337* \ingroup CUDA_MATH__BFLOAT16_MISC
338* \brief Converts high 16 bits of \p nv_bfloat162 to float and returns the result
339*
340* \details Converts high 16 bits of \p nv_bfloat162 input \p a to 32-bit floating-point number
341* and returns the result.
342* \param[in] a - nv_bfloat162. Is only being read.
343*
344* \returns float
345* - The high 16 bits of \p a converted to float.
346* \internal
347* \exception-guarantee no-throw guarantee
348* \behavior reentrant, thread safe
349* \endinternal
350*/
351__CUDA_HOSTDEVICE_BF16_DECL__ float __high2float(const __nv_bfloat162 a);
352
353/**
354* \ingroup CUDA_MATH__BFLOAT16_MISC
355* \brief Converts both components of float2 number to nv_bfloat16 precision in
356* round-to-nearest-even mode and returns \p nv_bfloat162 with converted values.
357*
358* \details Converts both components of float2 to nv_bfloat16 precision in round-to-nearest-even
359* mode and combines the results into one \p nv_bfloat162 number. Low 16 bits of the
360* return value correspond to \p a.x and high 16 bits of the return value
361* correspond to \p a.y.
362* \param[in] a - float2. Is only being read.
363*
364* \returns nv_bfloat162
365* - The \p nv_bfloat162 which has corresponding halves equal to the
366* converted float2 components.
367* \internal
368* \exception-guarantee no-throw guarantee
369* \behavior reentrant, thread safe
370* \endinternal
371*/
372__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat162 __float22bfloat162_rn(const float2 a);
373/**
374* \ingroup CUDA_MATH__BFLOAT16_MISC
375* \brief Converts both halves of \p nv_bfloat162 to float2 and returns the result.
376*
377* \details Converts both halves of \p nv_bfloat162 input \p a to float2 and returns the
378* result.
379* \param[in] a - nv_bfloat162. Is only being read.
380*
381* \returns float2
382* - \p a converted to float2.
383* \internal
384* \exception-guarantee no-throw guarantee
385* \behavior reentrant, thread safe
386* \endinternal
387*/
388__CUDA_HOSTDEVICE_BF16_DECL__ float2 __bfloat1622float2(const __nv_bfloat162 a);
389/**
390* \ingroup CUDA_MATH__BFLOAT16_MISC
391* \brief Convert a nv_bfloat16 to a signed char in round-towards-zero mode.
392*
393* \details Convert the nv_bfloat16 floating-point value \p h to a signed
394* char in round-towards-zero mode. NaN inputs are converted to 0.
395* \param[in] h - nv_bfloat16. Is only being read.
396*
397* \returns signed char
398* - \p h converted to a signed char.
399* \internal
400* \exception-guarantee no-throw guarantee
401* \behavior reentrant, thread safe
402* \endinternal
403*/
404__CUDA_HOSTDEVICE_BF16_DECL__ signed char __bfloat162char_rz(const __nv_bfloat16 h);
405/**
406* \ingroup CUDA_MATH__BFLOAT16_MISC
407* \brief Convert a nv_bfloat16 to an unsigned char in round-towards-zero mode.
408*
409* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned
410* char in round-towards-zero mode. NaN inputs are converted to 0.
411* \param[in] h - nv_bfloat16. Is only being read.
412*
413* \returns unsigned char
414* - \p h converted to an unsigned char.
415* \internal
416* \exception-guarantee no-throw guarantee
417* \behavior reentrant, thread safe
418* \endinternal
419*/
420__CUDA_HOSTDEVICE_BF16_DECL__ unsigned char __bfloat162uchar_rz(const __nv_bfloat16 h);
421#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
422/**
423* \ingroup CUDA_MATH__BFLOAT16_MISC
424* \brief Convert a nv_bfloat16 to a signed integer in round-to-nearest-even mode.
425*
426* \details Convert the nv_bfloat16 floating-point value \p h to a signed integer in
427* round-to-nearest-even mode. NaN inputs are converted to 0.
428* \param[in] h - nv_bfloat16. Is only being read.
429*
430* \returns int
431* - \p h converted to a signed integer.
432* \internal
433* \exception-guarantee no-throw guarantee
434* \behavior reentrant, thread safe
435* \endinternal
436*/
437__CUDA_BF16_DECL__ int __bfloat162int_rn(const __nv_bfloat16 h);
438#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
439/**
440* \ingroup CUDA_MATH__BFLOAT16_MISC
441* \brief Convert a nv_bfloat16 to a signed integer in round-towards-zero mode.
442*
443* \details Convert the nv_bfloat16 floating-point value \p h to a signed integer in
444* round-towards-zero mode. NaN inputs are converted to 0.
445* \param[in] h - nv_bfloat16. Is only being read.
446*
447* \returns int
448* - \p h converted to a signed integer.
449* \internal
450* \exception-guarantee no-throw guarantee
451* \behavior reentrant, thread safe
452* \endinternal
453*/
454__CUDA_HOSTDEVICE_BF16_DECL__ int __bfloat162int_rz(const __nv_bfloat16 h);
455#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
456/**
457* \ingroup CUDA_MATH__BFLOAT16_MISC
458* \brief Convert a nv_bfloat16 to a signed integer in round-down mode.
459*
460* \details Convert the nv_bfloat16 floating-point value \p h to a signed integer in
461* round-down mode. NaN inputs are converted to 0.
462* \param[in] h - nv_bfloat16. Is only being read.
463*
464* \returns int
465* - \p h converted to a signed integer.
466* \internal
467* \exception-guarantee no-throw guarantee
468* \behavior reentrant, thread safe
469* \endinternal
470*/
471__CUDA_BF16_DECL__ int __bfloat162int_rd(const __nv_bfloat16 h);
472/**
473* \ingroup CUDA_MATH__BFLOAT16_MISC
474* \brief Convert a nv_bfloat16 to a signed integer in round-up mode.
475*
476* \details Convert the nv_bfloat16 floating-point value \p h to a signed integer in
477* round-up mode. NaN inputs are converted to 0.
478* \param[in] h - nv_bfloat16. Is only being read.
479*
480* \returns int
481* - \p h converted to a signed integer.
482* \internal
483* \exception-guarantee no-throw guarantee
484* \behavior reentrant, thread safe
485* \endinternal
486*/
487__CUDA_BF16_DECL__ int __bfloat162int_ru(const __nv_bfloat16 h);
488#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
489/**
490* \ingroup CUDA_MATH__BFLOAT16_MISC
491* \brief Convert a signed integer to a nv_bfloat16 in round-to-nearest-even mode.
492*
493* \details Convert the signed integer value \p i to a nv_bfloat16 floating-point
494* value in round-to-nearest-even mode.
495* \param[in] i - int. Is only being read.
496*
497* \returns nv_bfloat16
498* - \p i converted to nv_bfloat16.
499* \internal
500* \exception-guarantee no-throw guarantee
501* \behavior reentrant, thread safe
502* \endinternal
503*/
504__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __int2bfloat16_rn(const int i);
505#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
506/**
507* \ingroup CUDA_MATH__BFLOAT16_MISC
508* \brief Convert a signed integer to a nv_bfloat16 in round-towards-zero mode.
509*
510* \details Convert the signed integer value \p i to a nv_bfloat16 floating-point
511* value in round-towards-zero mode.
512* \param[in] i - int. Is only being read.
513*
514* \returns nv_bfloat16
515* - \p i converted to nv_bfloat16.
516* \internal
517* \exception-guarantee no-throw guarantee
518* \behavior reentrant, thread safe
519* \endinternal
520*/
521__CUDA_BF16_DECL__ __nv_bfloat16 __int2bfloat16_rz(const int i);
522/**
523* \ingroup CUDA_MATH__BFLOAT16_MISC
524* \brief Convert a signed integer to a nv_bfloat16 in round-down mode.
525*
526* \details Convert the signed integer value \p i to a nv_bfloat16 floating-point
527* value in round-down mode.
528* \param[in] i - int. Is only being read.
529*
530* \returns nv_bfloat16
531* - \p i converted to nv_bfloat16.
532* \internal
533* \exception-guarantee no-throw guarantee
534* \behavior reentrant, thread safe
535* \endinternal
536*/
537__CUDA_BF16_DECL__ __nv_bfloat16 __int2bfloat16_rd(const int i);
538/**
539* \ingroup CUDA_MATH__BFLOAT16_MISC
540* \brief Convert a signed integer to a nv_bfloat16 in round-up mode.
541*
542* \details Convert the signed integer value \p i to a nv_bfloat16 floating-point
543* value in round-up mode.
544* \param[in] i - int. Is only being read.
545*
546* \returns nv_bfloat16
547* - \p i converted to nv_bfloat16.
548* \internal
549* \exception-guarantee no-throw guarantee
550* \behavior reentrant, thread safe
551* \endinternal
552*/
553__CUDA_BF16_DECL__ __nv_bfloat16 __int2bfloat16_ru(const int i);
554
555/**
556* \ingroup CUDA_MATH__BFLOAT16_MISC
557* \brief Convert a nv_bfloat16 to a signed short integer in round-to-nearest-even
558* mode.
559*
560* \details Convert the nv_bfloat16 floating-point value \p h to a signed short
561* integer in round-to-nearest-even mode. NaN inputs are converted to 0.
562* \param[in] h - nv_bfloat16. Is only being read.
563*
564* \returns short int
565* - \p h converted to a signed short integer.
566* \internal
567* \exception-guarantee no-throw guarantee
568* \behavior reentrant, thread safe
569* \endinternal
570*/
571__CUDA_BF16_DECL__ short int __bfloat162short_rn(const __nv_bfloat16 h);
572#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
573/**
574* \ingroup CUDA_MATH__BFLOAT16_MISC
575* \brief Convert a nv_bfloat16 to a signed short integer in round-towards-zero mode.
576*
577* \details Convert the nv_bfloat16 floating-point value \p h to a signed short
578* integer in round-towards-zero mode. NaN inputs are converted to 0.
579* \param[in] h - nv_bfloat16. Is only being read.
580*
581* \returns short int
582* - \p h converted to a signed short integer.
583* \internal
584* \exception-guarantee no-throw guarantee
585* \behavior reentrant, thread safe
586* \endinternal
587*/
588__CUDA_HOSTDEVICE_BF16_DECL__ short int __bfloat162short_rz(const __nv_bfloat16 h);
589#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
590/**
591* \ingroup CUDA_MATH__BFLOAT16_MISC
592* \brief Convert a nv_bfloat16 to a signed short integer in round-down mode.
593*
594* \details Convert the nv_bfloat16 floating-point value \p h to a signed short
595* integer in round-down mode. NaN inputs are converted to 0.
596* \param[in] h - nv_bfloat16. Is only being read.
597*
598* \returns short int
599* - \p h converted to a signed short integer.
600* \internal
601* \exception-guarantee no-throw guarantee
602* \behavior reentrant, thread safe
603* \endinternal
604*/
605__CUDA_BF16_DECL__ short int __bfloat162short_rd(const __nv_bfloat16 h);
606/**
607* \ingroup CUDA_MATH__BFLOAT16_MISC
608* \brief Convert a nv_bfloat16 to a signed short integer in round-up mode.
609*
610* \details Convert the nv_bfloat16 floating-point value \p h to a signed short
611* integer in round-up mode. NaN inputs are converted to 0.
612* \param[in] h - nv_bfloat16. Is only being read.
613*
614* \returns short int
615* - \p h converted to a signed short integer.
616* \internal
617* \exception-guarantee no-throw guarantee
618* \behavior reentrant, thread safe
619* \endinternal
620*/
621__CUDA_BF16_DECL__ short int __bfloat162short_ru(const __nv_bfloat16 h);
622#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
623/**
624* \ingroup CUDA_MATH__BFLOAT16_MISC
625* \brief Convert a signed short integer to a nv_bfloat16 in round-to-nearest-even
626* mode.
627*
628* \details Convert the signed short integer value \p i to a nv_bfloat16 floating-point
629* value in round-to-nearest-even mode.
630* \param[in] i - short int. Is only being read.
631*
632* \returns nv_bfloat16
633* - \p i converted to nv_bfloat16.
634* \internal
635* \exception-guarantee no-throw guarantee
636* \behavior reentrant, thread safe
637* \endinternal
638*/
639__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __short2bfloat16_rn(const short int i);
640#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
641/**
642* \ingroup CUDA_MATH__BFLOAT16_MISC
643* \brief Convert a signed short integer to a nv_bfloat16 in round-towards-zero mode.
644*
645* \details Convert the signed short integer value \p i to a nv_bfloat16 floating-point
646* value in round-towards-zero mode.
647* \param[in] i - short int. Is only being read.
648*
649* \returns nv_bfloat16
650* - \p i converted to nv_bfloat16.
651* \internal
652* \exception-guarantee no-throw guarantee
653* \behavior reentrant, thread safe
654* \endinternal
655*/
656__CUDA_BF16_DECL__ __nv_bfloat16 __short2bfloat16_rz(const short int i);
657/**
658* \ingroup CUDA_MATH__BFLOAT16_MISC
659* \brief Convert a signed short integer to a nv_bfloat16 in round-down mode.
660*
661* \details Convert the signed short integer value \p i to a nv_bfloat16 floating-point
662* value in round-down mode.
663* \param[in] i - short int. Is only being read.
664*
665* \returns nv_bfloat16
666* - \p i converted to nv_bfloat16.
667* \internal
668* \exception-guarantee no-throw guarantee
669* \behavior reentrant, thread safe
670* \endinternal
671*/
672__CUDA_BF16_DECL__ __nv_bfloat16 __short2bfloat16_rd(const short int i);
673/**
674* \ingroup CUDA_MATH__BFLOAT16_MISC
675* \brief Convert a signed short integer to a nv_bfloat16 in round-up mode.
676*
677* \details Convert the signed short integer value \p i to a nv_bfloat16 floating-point
678* value in round-up mode.
679* \param[in] i - short int. Is only being read.
680*
681* \returns nv_bfloat16
682* - \p i converted to nv_bfloat16.
683* \internal
684* \exception-guarantee no-throw guarantee
685* \behavior reentrant, thread safe
686* \endinternal
687*/
688__CUDA_BF16_DECL__ __nv_bfloat16 __short2bfloat16_ru(const short int i);
689/**
690* \ingroup CUDA_MATH__BFLOAT16_MISC
691* \brief Convert a nv_bfloat16 to an unsigned integer in round-to-nearest-even mode.
692*
693* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned integer
694* in round-to-nearest-even mode. NaN inputs are converted to 0.
695* \param[in] h - nv_bfloat16. Is only being read.
696*
697* \returns unsigned int
698* - \p h converted to an unsigned integer.
699* \internal
700* \exception-guarantee no-throw guarantee
701* \behavior reentrant, thread safe
702* \endinternal
703*/
704__CUDA_BF16_DECL__ unsigned int __bfloat162uint_rn(const __nv_bfloat16 h);
705#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
706
707/**
708* \ingroup CUDA_MATH__BFLOAT16_MISC
709* \brief Convert a nv_bfloat16 to an unsigned integer in round-towards-zero mode.
710*
711* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned integer
712* in round-towards-zero mode. NaN inputs are converted to 0.
713* \param[in] h - nv_bfloat16. Is only being read.
714*
715* \returns unsigned int
716* - \p h converted to an unsigned integer.
717* \internal
718* \exception-guarantee no-throw guarantee
719* \behavior reentrant, thread safe
720* \endinternal
721*/
722__CUDA_HOSTDEVICE_BF16_DECL__ unsigned int __bfloat162uint_rz(const __nv_bfloat16 h);
723#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
724/**
725* \ingroup CUDA_MATH__BFLOAT16_MISC
726* \brief Convert a nv_bfloat16 to an unsigned integer in round-down mode.
727*
728* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned integer
729* in round-down mode. NaN inputs are converted to 0.
730* \param[in] h - nv_bfloat16. Is only being read.
731*
732* \returns unsigned int
733* - \p h converted to an unsigned integer.
734* \internal
735* \exception-guarantee no-throw guarantee
736* \behavior reentrant, thread safe
737* \endinternal
738*/
739__CUDA_BF16_DECL__ unsigned int __bfloat162uint_rd(const __nv_bfloat16 h);
740/**
741* \ingroup CUDA_MATH__BFLOAT16_MISC
742* \brief Convert a nv_bfloat16 to an unsigned integer in round-up mode.
743*
744* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned integer
745* in round-up mode. NaN inputs are converted to 0.
746* \param[in] h - nv_bfloat16. Is only being read.
747*
748* \returns unsigned int
749* - \p h converted to an unsigned integer.
750* \internal
751* \exception-guarantee no-throw guarantee
752* \behavior reentrant, thread safe
753* \endinternal
754*/
755__CUDA_BF16_DECL__ unsigned int __bfloat162uint_ru(const __nv_bfloat16 h);
756#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
757
758/**
759* \ingroup CUDA_MATH__BFLOAT16_MISC
760* \brief Convert an unsigned integer to a nv_bfloat16 in round-to-nearest-even mode.
761*
762* \details Convert the unsigned integer value \p i to a nv_bfloat16 floating-point
763* value in round-to-nearest-even mode.
764* \param[in] i - unsigned int. Is only being read.
765*
766* \returns nv_bfloat16
767* - \p i converted to nv_bfloat16.
768* \internal
769* \exception-guarantee no-throw guarantee
770* \behavior reentrant, thread safe
771* \endinternal
772*/
773__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __uint2bfloat16_rn(const unsigned int i);
774#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
775/**
776* \ingroup CUDA_MATH__BFLOAT16_MISC
777* \brief Convert an unsigned integer to a nv_bfloat16 in round-towards-zero mode.
778*
779* \details Convert the unsigned integer value \p i to a nv_bfloat16 floating-point
780* value in round-towards-zero mode.
781* \param[in] i - unsigned int. Is only being read.
782*
783* \returns nv_bfloat16
784* - \p i converted to nv_bfloat16.
785* \internal
786* \exception-guarantee no-throw guarantee
787* \behavior reentrant, thread safe
788* \endinternal
789*/
790__CUDA_BF16_DECL__ __nv_bfloat16 __uint2bfloat16_rz(const unsigned int i);
791/**
792* \ingroup CUDA_MATH__BFLOAT16_MISC
793* \brief Convert an unsigned integer to a nv_bfloat16 in round-down mode.
794*
795* \details Convert the unsigned integer value \p i to a nv_bfloat16 floating-point
796* value in round-down mode.
797* \param[in] i - unsigned int. Is only being read.
798*
799* \returns nv_bfloat16
800* - \p i converted to nv_bfloat16.
801* \internal
802* \exception-guarantee no-throw guarantee
803* \behavior reentrant, thread safe
804* \endinternal
805*/
806__CUDA_BF16_DECL__ __nv_bfloat16 __uint2bfloat16_rd(const unsigned int i);
807/**
808* \ingroup CUDA_MATH__BFLOAT16_MISC
809* \brief Convert an unsigned integer to a nv_bfloat16 in round-up mode.
810*
811* \details Convert the unsigned integer value \p i to a nv_bfloat16 floating-point
812* value in round-up mode.
813* \param[in] i - unsigned int. Is only being read.
814*
815* \returns nv_bfloat16
816* - \p i converted to nv_bfloat16.
817* \internal
818* \exception-guarantee no-throw guarantee
819* \behavior reentrant, thread safe
820* \endinternal
821*/
822__CUDA_BF16_DECL__ __nv_bfloat16 __uint2bfloat16_ru(const unsigned int i);
823
824/**
825* \ingroup CUDA_MATH__BFLOAT16_MISC
826* \brief Convert a nv_bfloat16 to an unsigned short integer in round-to-nearest-even
827* mode.
828*
829* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned short
830* integer in round-to-nearest-even mode. NaN inputs are converted to 0.
831* \param[in] h - nv_bfloat16. Is only being read.
832*
833* \returns unsigned short int
834* - \p h converted to an unsigned short integer.
835* \internal
836* \exception-guarantee no-throw guarantee
837* \behavior reentrant, thread safe
838* \endinternal
839*/
840__CUDA_BF16_DECL__ unsigned short int __bfloat162ushort_rn(const __nv_bfloat16 h);
841#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
842
843/**
844* \ingroup CUDA_MATH__BFLOAT16_MISC
845* \brief Convert a nv_bfloat16 to an unsigned short integer in round-towards-zero
846* mode.
847*
848* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned short
849* integer in round-towards-zero mode. NaN inputs are converted to 0.
850* \param[in] h - nv_bfloat16. Is only being read.
851*
852* \returns unsigned short int
853* - \p h converted to an unsigned short integer.
854* \internal
855* \exception-guarantee no-throw guarantee
856* \behavior reentrant, thread safe
857* \endinternal
858*/
859__CUDA_HOSTDEVICE_BF16_DECL__ unsigned short int __bfloat162ushort_rz(const __nv_bfloat16 h);
860#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
861/**
862* \ingroup CUDA_MATH__BFLOAT16_MISC
863* \brief Convert a nv_bfloat16 to an unsigned short integer in round-down mode.
864*
865* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned short
866* integer in round-down mode. NaN inputs are converted to 0.
867* \param[in] h - nv_bfloat16. Is only being read.
868*
869* \returns unsigned short int
870* - \p h converted to an unsigned short integer.
871*/
872__CUDA_BF16_DECL__ unsigned short int __bfloat162ushort_rd(const __nv_bfloat16 h);
873/**
874* \ingroup CUDA_MATH__BFLOAT16_MISC
875* \brief Convert a nv_bfloat16 to an unsigned short integer in round-up mode.
876*
877* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned short
878* integer in round-up mode. NaN inputs are converted to 0.
879* \param[in] h - nv_bfloat16. Is only being read.
880*
881* \returns unsigned short int
882* - \p h converted to an unsigned short integer.
883*/
884__CUDA_BF16_DECL__ unsigned short int __bfloat162ushort_ru(const __nv_bfloat16 h);
885#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
886
887/**
888* \ingroup CUDA_MATH__BFLOAT16_MISC
889* \brief Convert an unsigned short integer to a nv_bfloat16 in round-to-nearest-even
890* mode.
891*
892* \details Convert the unsigned short integer value \p i to a nv_bfloat16 floating-point
893* value in round-to-nearest-even mode.
894* \param[in] i - unsigned short int. Is only being read.
895*
896* \returns nv_bfloat16
897* - \p i converted to nv_bfloat16.
898* \internal
899* \exception-guarantee no-throw guarantee
900* \behavior reentrant, thread safe
901* \endinternal
902*/
903__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __ushort2bfloat16_rn(const unsigned short int i);
904#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
905/**
906* \ingroup CUDA_MATH__BFLOAT16_MISC
907* \brief Convert an unsigned short integer to a nv_bfloat16 in round-towards-zero
908* mode.
909*
910* \details Convert the unsigned short integer value \p i to a nv_bfloat16 floating-point
911* value in round-towards-zero mode.
912* \param[in] i - unsigned short int. Is only being read.
913*
914* \returns nv_bfloat16
915* - \p i converted to nv_bfloat16.
916* \internal
917* \exception-guarantee no-throw guarantee
918* \behavior reentrant, thread safe
919* \endinternal
920*/
921__CUDA_BF16_DECL__ __nv_bfloat16 __ushort2bfloat16_rz(const unsigned short int i);
922/**
923* \ingroup CUDA_MATH__BFLOAT16_MISC
924* \brief Convert an unsigned short integer to a nv_bfloat16 in round-down mode.
925*
926* \details Convert the unsigned short integer value \p i to a nv_bfloat16 floating-point
927* value in round-down mode.
928* \param[in] i - unsigned short int. Is only being read.
929*
930* \returns nv_bfloat16
931* - \p i converted to nv_bfloat16.
932* \internal
933* \exception-guarantee no-throw guarantee
934* \behavior reentrant, thread safe
935* \endinternal
936*/
937__CUDA_BF16_DECL__ __nv_bfloat16 __ushort2bfloat16_rd(const unsigned short int i);
938/**
939* \ingroup CUDA_MATH__BFLOAT16_MISC
940* \brief Convert an unsigned short integer to a nv_bfloat16 in round-up mode.
941*
942* \details Convert the unsigned short integer value \p i to a nv_bfloat16 floating-point
943* value in round-up mode.
944* \param[in] i - unsigned short int. Is only being read.
945*
946* \returns nv_bfloat16
947* - \p i converted to nv_bfloat16.
948* \internal
949* \exception-guarantee no-throw guarantee
950* \behavior reentrant, thread safe
951* \endinternal
952*/
953__CUDA_BF16_DECL__ __nv_bfloat16 __ushort2bfloat16_ru(const unsigned short int i);
954/**
955* \ingroup CUDA_MATH__BFLOAT16_MISC
956* \brief Convert a nv_bfloat16 to an unsigned 64-bit integer in round-to-nearest-even
957* mode.
958*
959* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned 64-bit
960* integer in round-to-nearest-even mode. NaN inputs return 0x8000000000000000.
961* \param[in] h - nv_bfloat16. Is only being read.
962*
963* \returns unsigned long long int
964* - \p h converted to an unsigned 64-bit integer.
965* \internal
966* \exception-guarantee no-throw guarantee
967* \behavior reentrant, thread safe
968* \endinternal
969*/
970__CUDA_BF16_DECL__ unsigned long long int __bfloat162ull_rn(const __nv_bfloat16 h);
971#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
972
973/**
974* \ingroup CUDA_MATH__BFLOAT16_MISC
975* \brief Convert a nv_bfloat16 to an unsigned 64-bit integer in round-towards-zero
976* mode.
977*
978* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned 64-bit
979* integer in round-towards-zero mode. NaN inputs return 0x8000000000000000.
980* \param[in] h - nv_bfloat16. Is only being read.
981*
982* \returns unsigned long long int
983* - \p h converted to an unsigned 64-bit integer.
984* \internal
985* \exception-guarantee no-throw guarantee
986* \behavior reentrant, thread safe
987* \endinternal
988*/
989__CUDA_HOSTDEVICE_BF16_DECL__ unsigned long long int __bfloat162ull_rz(const __nv_bfloat16 h);
990/**
991* \ingroup CUDA_MATH__BFLOAT16_MISC
992* \brief Vector function, combines two \p nv_bfloat16 numbers into one \p nv_bfloat162 number.
993*
994* \details Combines two input \p nv_bfloat16 number \p x and \p y into one \p nv_bfloat162 number.
995* Input \p x is stored in low 16 bits of the return value, input \p y is stored
996* in high 16 bits of the return value.
997* \param[in] x - nv_bfloat16. Is only being read.
998* \param[in] y - nv_bfloat16. Is only being read.
999*
1000* \returns __nv_bfloat162
1001* - The \p __nv_bfloat162 vector with one half equal to \p x and the other to \p y.
1002* \internal
1003* \exception-guarantee no-throw guarantee
1004* \behavior reentrant, thread safe
1005* \endinternal
1006*/
1007__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat162 make_bfloat162(const __nv_bfloat16 x, const __nv_bfloat16 y);
1008#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
1009/**
1010* \ingroup CUDA_MATH__BFLOAT16_MISC
1011* \brief Convert a nv_bfloat16 to an unsigned 64-bit integer in round-down mode.
1012*
1013* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned 64-bit
1014* integer in round-down mode. NaN inputs return 0x8000000000000000.
1015* \param[in] h - nv_bfloat16. Is only being read.
1016*
1017* \returns unsigned long long int
1018* - \p h converted to an unsigned 64-bit integer.
1019* \internal
1020* \exception-guarantee no-throw guarantee
1021* \behavior reentrant, thread safe
1022* \endinternal
1023*/
1024__CUDA_BF16_DECL__ unsigned long long int __bfloat162ull_rd(const __nv_bfloat16 h);
1025/**
1026* \ingroup CUDA_MATH__BFLOAT16_MISC
1027* \brief Convert a nv_bfloat16 to an unsigned 64-bit integer in round-up mode.
1028*
1029* \details Convert the nv_bfloat16 floating-point value \p h to an unsigned 64-bit
1030* integer in round-up mode. NaN inputs return 0x8000000000000000.
1031* \param[in] h - nv_bfloat16. Is only being read.
1032*
1033* \returns unsigned long long int
1034* - \p h converted to an unsigned 64-bit integer.
1035* \internal
1036* \exception-guarantee no-throw guarantee
1037* \behavior reentrant, thread safe
1038* \endinternal
1039*/
1040__CUDA_BF16_DECL__ unsigned long long int __bfloat162ull_ru(const __nv_bfloat16 h);
1041#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
1042/**
1043* \ingroup CUDA_MATH__BFLOAT16_MISC
1044* \brief Convert an unsigned 64-bit integer to a nv_bfloat16 in round-to-nearest-even
1045* mode.
1046*
1047* \details Convert the unsigned 64-bit integer value \p i to a nv_bfloat16 floating-point
1048* value in round-to-nearest-even mode.
1049* \param[in] i - unsigned long long int. Is only being read.
1050*
1051* \returns nv_bfloat16
1052* - \p i converted to nv_bfloat16.
1053* \internal
1054* \exception-guarantee no-throw guarantee
1055* \behavior reentrant, thread safe
1056* \endinternal
1057*/
1058__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __ull2bfloat16_rn(const unsigned long long int i);
1059#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
1060/**
1061* \ingroup CUDA_MATH__BFLOAT16_MISC
1062* \brief Convert an unsigned 64-bit integer to a nv_bfloat16 in round-towards-zero
1063* mode.
1064*
1065* \details Convert the unsigned 64-bit integer value \p i to a nv_bfloat16 floating-point
1066* value in round-towards-zero mode.
1067* \param[in] i - unsigned long long int. Is only being read.
1068*
1069* \returns nv_bfloat16
1070* - \p i converted to nv_bfloat16.
1071* \internal
1072* \exception-guarantee no-throw guarantee
1073* \behavior reentrant, thread safe
1074* \endinternal
1075*/
1076__CUDA_BF16_DECL__ __nv_bfloat16 __ull2bfloat16_rz(const unsigned long long int i);
1077/**
1078* \ingroup CUDA_MATH__BFLOAT16_MISC
1079* \brief Convert an unsigned 64-bit integer to a nv_bfloat16 in round-down mode.
1080*
1081* \details Convert the unsigned 64-bit integer value \p i to a nv_bfloat16 floating-point
1082* value in round-down mode.
1083* \param[in] i - unsigned long long int. Is only being read.
1084*
1085* \returns nv_bfloat16
1086* - \p i converted to nv_bfloat16.
1087* \internal
1088* \exception-guarantee no-throw guarantee
1089* \behavior reentrant, thread safe
1090* \endinternal
1091*/
1092__CUDA_BF16_DECL__ __nv_bfloat16 __ull2bfloat16_rd(const unsigned long long int i);
1093/**
1094* \ingroup CUDA_MATH__BFLOAT16_MISC
1095* \brief Convert an unsigned 64-bit integer to a nv_bfloat16 in round-up mode.
1096*
1097* \details Convert the unsigned 64-bit integer value \p i to a nv_bfloat16 floating-point
1098* value in round-up mode.
1099* \param[in] i - unsigned long long int. Is only being read.
1100*
1101* \returns nv_bfloat16
1102* - \p i converted to nv_bfloat16.
1103* \internal
1104* \exception-guarantee no-throw guarantee
1105* \behavior reentrant, thread safe
1106* \endinternal
1107*/
1108__CUDA_BF16_DECL__ __nv_bfloat16 __ull2bfloat16_ru(const unsigned long long int i);
1109/**
1110* \ingroup CUDA_MATH__BFLOAT16_MISC
1111* \brief Convert a nv_bfloat16 to a signed 64-bit integer in round-to-nearest-even
1112* mode.
1113*
1114* \details Convert the nv_bfloat16 floating-point value \p h to a signed 64-bit
1115* integer in round-to-nearest-even mode. NaN inputs return a long long int with hex value of 0x8000000000000000.
1116* \param[in] h - nv_bfloat16. Is only being read.
1117*
1118* \returns long long int
1119* - \p h converted to a signed 64-bit integer.
1120* \internal
1121* \exception-guarantee no-throw guarantee
1122* \behavior reentrant, thread safe
1123* \endinternal
1124*/
1125__CUDA_BF16_DECL__ long long int __bfloat162ll_rn(const __nv_bfloat16 h);
1126#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
1127/**
1128* \ingroup CUDA_MATH__BFLOAT16_MISC
1129* \brief Convert a nv_bfloat16 to a signed 64-bit integer in round-towards-zero mode.
1130*
1131* \details Convert the nv_bfloat16 floating-point value \p h to a signed 64-bit
1132* integer in round-towards-zero mode. NaN inputs return a long long int with hex value of 0x8000000000000000.
1133* \param[in] h - nv_bfloat16. Is only being read.
1134*
1135* \returns long long int
1136* - \p h converted to a signed 64-bit integer.
1137* \internal
1138* \exception-guarantee no-throw guarantee
1139* \behavior reentrant, thread safe
1140* \endinternal
1141*/
1142__CUDA_HOSTDEVICE_BF16_DECL__ long long int __bfloat162ll_rz(const __nv_bfloat16 h);
1143#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
1144/**
1145* \ingroup CUDA_MATH__BFLOAT16_MISC
1146* \brief Convert a nv_bfloat16 to a signed 64-bit integer in round-down mode.
1147*
1148* \details Convert the nv_bfloat16 floating-point value \p h to a signed 64-bit
1149* integer in round-down mode. NaN inputs return a long long int with hex value of 0x8000000000000000.
1150* \param[in] h - nv_bfloat16. Is only being read.
1151*
1152* \returns long long int
1153* - \p h converted to a signed 64-bit integer.
1154* \internal
1155* \exception-guarantee no-throw guarantee
1156* \behavior reentrant, thread safe
1157* \endinternal
1158*/
1159__CUDA_BF16_DECL__ long long int __bfloat162ll_rd(const __nv_bfloat16 h);
1160/**
1161* \ingroup CUDA_MATH__BFLOAT16_MISC
1162* \brief Convert a nv_bfloat16 to a signed 64-bit integer in round-up mode.
1163*
1164* \details Convert the nv_bfloat16 floating-point value \p h to a signed 64-bit
1165* integer in round-up mode. NaN inputs return a long long int with hex value of 0x8000000000000000.
1166* \param[in] h - nv_bfloat16. Is only being read.
1167*
1168* \returns long long int
1169* - \p h converted to a signed 64-bit integer.
1170* \internal
1171* \exception-guarantee no-throw guarantee
1172* \behavior reentrant, thread safe
1173* \endinternal
1174*/
1175__CUDA_BF16_DECL__ long long int __bfloat162ll_ru(const __nv_bfloat16 h);
1176#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
1177/**
1178* \ingroup CUDA_MATH__BFLOAT16_MISC
1179* \brief Convert a signed 64-bit integer to a nv_bfloat16 in round-to-nearest-even
1180* mode.
1181*
1182* \details Convert the signed 64-bit integer value \p i to a nv_bfloat16 floating-point
1183* value in round-to-nearest-even mode.
1184* \param[in] i - long long int. Is only being read.
1185*
1186* \returns nv_bfloat16
1187* - \p i converted to nv_bfloat16.
1188* \internal
1189* \exception-guarantee no-throw guarantee
1190* \behavior reentrant, thread safe
1191* \endinternal
1192*/
1193__CUDA_HOSTDEVICE_BF16_DECL__ __nv_bfloat16 __ll2bfloat16_rn(const long long int i);
1194#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
1195/**
1196* \ingroup CUDA_MATH__BFLOAT16_MISC
1197* \brief Convert a signed 64-bit integer to a nv_bfloat16 in round-towards-zero mode.
1198*
1199* \details Convert the signed 64-bit integer value \p i to a nv_bfloat16 floating-point
1200* value in round-towards-zero mode.