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_HALF Half Precision Intrinsics
52* This section describes half precision intrinsic functions.
53* To use these functions, include the header file \p cuda_fp16.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.
61*
62* The following macros are available to help users selectively enable/disable
63* various definitions present in the header file:
64* - \p CUDA_NO_HALF - If defined, this macro will prevent the definition of
65* additional type aliases in the global namespace, helping to avoid potential
66* conflicts with symbols defined in the user program.
67* - \p __CUDA_NO_HALF_CONVERSIONS__ - If defined, this macro will prevent the
68* use of the C++ type conversions (converting constructors and conversion
69* operators) that are common for built-in floating-point types, but may be
70* undesirable for \p half which is essentially a user-defined type.
71* - \p __CUDA_NO_HALF_OPERATORS__ and \p __CUDA_NO_HALF2_OPERATORS__ - If
72* defined, these macros will prevent the inadvertent use of usual arithmetic
73* and comparison operators. This enforces the storage-only type semantics and
74* prevents C++ style computations on \p half and \p half2 types.
75*/
76
77/**
78* \defgroup CUDA_MATH_INTRINSIC_HALF_CONSTANTS Half Arithmetic Constants
79* \ingroup CUDA_MATH_INTRINSIC_HALF
80* To use these constants, include the header file \p cuda_fp16.h in your program.
81*/
82
83/**
84* \defgroup CUDA_MATH__HALF_ARITHMETIC Half Arithmetic Functions
85* \ingroup CUDA_MATH_INTRINSIC_HALF
86* To use these functions, include the header file \p cuda_fp16.h in your program.
87*/
88
89/**
90* \defgroup CUDA_MATH__HALF2_ARITHMETIC Half2 Arithmetic Functions
91* \ingroup CUDA_MATH_INTRINSIC_HALF
92* To use these functions, include the header file \p cuda_fp16.h in your program.
93*/
94
95/**
96* \defgroup CUDA_MATH__HALF_COMPARISON Half Comparison Functions
97* \ingroup CUDA_MATH_INTRINSIC_HALF
98* To use these functions, include the header file \p cuda_fp16.h in your program.
99*/
100
101/**
102* \defgroup CUDA_MATH__HALF2_COMPARISON Half2 Comparison Functions
103* \ingroup CUDA_MATH_INTRINSIC_HALF
104* To use these functions, include the header file \p cuda_fp16.h in your program.
105*/
106
107/**
108* \defgroup CUDA_MATH__HALF_MISC Half Precision Conversion and Data Movement
109* \ingroup CUDA_MATH_INTRINSIC_HALF
110* To use these functions, include the header file \p cuda_fp16.h in your program.
111*/
112
113/**
114* \defgroup CUDA_MATH__HALF_FUNCTIONS Half Math Functions
115* \ingroup CUDA_MATH_INTRINSIC_HALF
116* To use these functions, include the header file \p cuda_fp16.h in your program.
117*/
118
119/**
120* \defgroup CUDA_MATH__HALF2_FUNCTIONS Half2 Math Functions
121* \ingroup CUDA_MATH_INTRINSIC_HALF
122* To use these functions, include the header file \p cuda_fp16.h in your program.
123*/
124
125#ifndef __CUDA_FP16_H__
126#define __CUDA_FP16_H__
127
128/* bring in float2, double4, etc vector types */
129#include "vector_types.h"
130/* bring in operations on vector types like: make_float2 */
131#include "vector_functions.h"
132
133#define ___CUDA_FP16_STRINGIFY_INNERMOST(x) #x
134#define __CUDA_FP16_STRINGIFY(x) ___CUDA_FP16_STRINGIFY_INNERMOST(x)
135
136#if defined(__cplusplus)
137
138/* Set up function decorations */
139#if (defined(__CUDACC_RTC__) && ((__CUDACC_VER_MAJOR__ > 12) || ((__CUDACC_VER_MAJOR__ == 12) && (__CUDACC_VER_MINOR__ >= 3))))
140#define __CUDA_FP16_DECL__ __device__
141#define __CUDA_HOSTDEVICE_FP16_DECL__ __device__
142#define __CUDA_HOSTDEVICE__ __device__
143#elif defined(__CUDACC__) || defined(_NVHPC_CUDA)
144#define __CUDA_FP16_DECL__ static __device__ __inline__
145#define __CUDA_HOSTDEVICE_FP16_DECL__ static __host__ __device__ __inline__
146#define __CUDA_HOSTDEVICE__ __host__ __device__
147#else /* !defined(__CUDACC__) */
148#if defined(__GNUC__)
149#define __CUDA_HOSTDEVICE_FP16_DECL__ static __attribute__ ((unused))
150#else
151#define __CUDA_HOSTDEVICE_FP16_DECL__ static
152#endif /* defined(__GNUC__) */
153#define __CUDA_HOSTDEVICE__
154#endif /* (defined(__CUDACC_RTC__) && ((__CUDACC_VER_MAJOR__ > 12) || ((__CUDACC_VER_MAJOR__ == 12) && (__CUDACC_VER_MINOR__ >= 3)))) */
155
156#define __CUDA_FP16_TYPES_EXIST__
157
158/* Macros to allow half & half2 to be used by inline assembly */
159#define __HALF_TO_US(var) *(reinterpret_cast<unsigned short *>(&(var)))
160#define __HALF_TO_CUS(var) *(reinterpret_cast<const unsigned short *>(&(var)))
161#define __HALF2_TO_UI(var) *(reinterpret_cast<unsigned int *>(&(var)))
162#define __HALF2_TO_CUI(var) *(reinterpret_cast<const unsigned int *>(&(var)))
163
164/* Forward-declaration of structures defined in "cuda_fp16.hpp" */
165struct __half;
166struct __half2;
167
168/**
169* \ingroup CUDA_MATH__HALF_MISC
170* \brief Converts double number to half precision in round-to-nearest-even mode
171* and returns \p half with converted value.
172*
173* \details Converts double number \p a to half precision in round-to-nearest-even mode.
174* \param[in] a - double. Is only being read.
175* \returns half
176* - \p a converted to half.
177* \internal
178* \exception-guarantee no-throw guarantee
179* \behavior reentrant, thread safe
180* \endinternal
181*/
182__CUDA_HOSTDEVICE_FP16_DECL__ __half __double2half(const double a);
183/**
184* \ingroup CUDA_MATH__HALF_MISC
185* \brief Converts float number to half precision in round-to-nearest-even mode
186* and returns \p half with converted value.
187*
188* \details Converts float number \p a to half precision in round-to-nearest-even mode.
189* \param[in] a - float. Is only being read.
190* \returns half
191* - \p a converted to half.
192* \internal
193* \exception-guarantee no-throw guarantee
194* \behavior reentrant, thread safe
195* \endinternal
196*/
197__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half(const float a);
198/**
199* \ingroup CUDA_MATH__HALF_MISC
200* \brief Converts float number to half precision in round-to-nearest-even mode
201* and returns \p half with converted value.
202*
203* \details Converts float number \p a to half precision in round-to-nearest-even mode.
204* \param[in] a - float. Is only being read.
205* \returns half
206* - \p a converted to half.
207* \internal
208* \exception-guarantee no-throw guarantee
209* \behavior reentrant, thread safe
210* \endinternal
211*/
212__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half_rn(const float a);
213/**
214* \ingroup CUDA_MATH__HALF_MISC
215* \brief Converts float number to half precision in round-towards-zero mode
216* and returns \p half with converted value.
217*
218* \details Converts float number \p a to half precision in round-towards-zero mode.
219* \param[in] a - float. Is only being read.
220* \returns half
221* - \p a converted to half.
222* \internal
223* \exception-guarantee no-throw guarantee
224* \behavior reentrant, thread safe
225* \endinternal
226*/
227__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half_rz(const float a);
228/**
229* \ingroup CUDA_MATH__HALF_MISC
230* \brief Converts float number to half precision in round-down mode
231* and returns \p half with converted value.
232*
233* \details Converts float number \p a to half precision in round-down mode.
234* \param[in] a - float. Is only being read.
235*
236* \returns half
237* - \p a converted to half.
238* \internal
239* \exception-guarantee no-throw guarantee
240* \behavior reentrant, thread safe
241* \endinternal
242*/
243__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half_rd(const float a);
244/**
245* \ingroup CUDA_MATH__HALF_MISC
246* \brief Converts float number to half precision in round-up mode
247* and returns \p half with converted value.
248*
249* \details Converts float number \p a to half precision in round-up mode.
250* \param[in] a - float. Is only being read.
251*
252* \returns half
253* - \p a converted to half.
254* \internal
255* \exception-guarantee no-throw guarantee
256* \behavior reentrant, thread safe
257* \endinternal
258*/
259__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half_ru(const float a);
260/**
261* \ingroup CUDA_MATH__HALF_MISC
262* \brief Converts \p half number to float.
263*
264* \details Converts half number \p a to float.
265* \param[in] a - float. Is only being read.
266*
267* \returns float
268* - \p a converted to float.
269* \internal
270* \exception-guarantee no-throw guarantee
271* \behavior reentrant, thread safe
272* \endinternal
273*/
274__CUDA_HOSTDEVICE_FP16_DECL__ float __half2float(const __half a);
275/**
276* \ingroup CUDA_MATH__HALF_MISC
277* \brief Converts input to half precision in round-to-nearest-even mode and
278* populates both halves of \p half2 with converted value.
279*
280* \details Converts input \p a to half precision in round-to-nearest-even mode and
281* populates both halves of \p half2 with converted value.
282* \param[in] a - float. Is only being read.
283*
284* \returns half2
285* - The \p half2 value with both halves equal to the converted half
286* precision number.
287* \internal
288* \exception-guarantee no-throw guarantee
289* \behavior reentrant, thread safe
290* \endinternal
291*/
292__CUDA_HOSTDEVICE_FP16_DECL__ __half2 __float2half2_rn(const float a);
293/**
294* \ingroup CUDA_MATH__HALF_MISC
295* \brief Converts both input floats to half precision in round-to-nearest-even
296* mode and returns \p half2 with converted values.
297*
298* \details Converts both input floats to half precision in round-to-nearest-even mode
299* and combines the results into one \p half2 number. Low 16 bits of the return
300* value correspond to the input \p a, high 16 bits correspond to the input \p
301* b.
302* \param[in] a - float. Is only being read.
303* \param[in] b - float. Is only being read.
304*
305* \returns half2
306* - The \p half2 value with corresponding halves equal to the
307* converted input floats.
308* \internal
309* \exception-guarantee no-throw guarantee
310* \behavior reentrant, thread safe
311* \endinternal
312*/
313__CUDA_HOSTDEVICE_FP16_DECL__ __half2 __floats2half2_rn(const float a, const float b);
314/**
315* \ingroup CUDA_MATH__HALF_MISC
316* \brief Converts low 16 bits of \p half2 to float and returns the result
317*
318* \details Converts low 16 bits of \p half2 input \p a to 32-bit floating-point number
319* and returns the result.
320* \param[in] a - half2. Is only being read.
321*
322* \returns float
323* - The low 16 bits of \p a converted to float.
324* \internal
325* \exception-guarantee no-throw guarantee
326* \behavior reentrant, thread safe
327* \endinternal
328*/
329__CUDA_HOSTDEVICE_FP16_DECL__ float __low2float(const __half2 a);
330/**
331* \ingroup CUDA_MATH__HALF_MISC
332* \brief Converts high 16 bits of \p half2 to float and returns the result
333*
334* \details Converts high 16 bits of \p half2 input \p a to 32-bit floating-point number
335* and returns the result.
336* \param[in] a - half2. Is only being read.
337*
338* \returns float
339* - The high 16 bits of \p a converted to float.
340* \internal
341* \exception-guarantee no-throw guarantee
342* \behavior reentrant, thread safe
343* \endinternal
344*/
345__CUDA_HOSTDEVICE_FP16_DECL__ float __high2float(const __half2 a);
346/**
347* \ingroup CUDA_MATH__HALF_MISC
348* \brief Convert a half to a signed char in round-towards-zero mode.
349*
350* \details Convert the half-precision floating-point value \p h to a signed char
351* integer in round-towards-zero mode. NaN inputs are converted to 0.
352* \param[in] h - half. Is only being read.
353*
354* \returns signed char
355* - \p h converted to a signed char.
356* \internal
357* \exception-guarantee no-throw guarantee
358* \behavior reentrant, thread safe
359* \endinternal
360*/
361__CUDA_HOSTDEVICE_FP16_DECL__ signed char __half2char_rz(const __half h);
362/**
363* \ingroup CUDA_MATH__HALF_MISC
364* \brief Convert a half to an unsigned char in round-towards-zero
365* mode.
366*
367* \details Convert the half-precision floating-point value \p h to an unsigned
368* char in round-towards-zero mode. NaN inputs are converted to 0.
369* \param[in] h - half. Is only being read.
370*
371* \returns unsigned char
372* - \p h converted to an unsigned char.
373* \internal
374* \exception-guarantee no-throw guarantee
375* \behavior reentrant, thread safe
376* \endinternal
377*/
378__CUDA_HOSTDEVICE_FP16_DECL__ unsigned char __half2uchar_rz(const __half h);
379/**
380* \ingroup CUDA_MATH__HALF_MISC
381* \brief Convert a half to a signed short integer in round-towards-zero mode.
382*
383* \details Convert the half-precision floating-point value \p h to a signed short
384* integer in round-towards-zero mode. NaN inputs are converted to 0.
385* \param[in] h - half. Is only being read.
386*
387* \returns short int
388* - \p h converted to a signed short integer.
389* \internal
390* \exception-guarantee no-throw guarantee
391* \behavior reentrant, thread safe
392* \endinternal
393*/
394__CUDA_HOSTDEVICE_FP16_DECL__ short int __half2short_rz(const __half h);
395/**
396* \ingroup CUDA_MATH__HALF_MISC
397* \brief Convert a half to an unsigned short integer in round-towards-zero
398* mode.
399*
400* \details Convert the half-precision floating-point value \p h to an unsigned short
401* integer in round-towards-zero mode. NaN inputs are converted to 0.
402* \param[in] h - half. Is only being read.
403*
404* \returns unsigned short int
405* - \p h converted to an unsigned short integer.
406* \internal
407* \exception-guarantee no-throw guarantee
408* \behavior reentrant, thread safe
409* \endinternal
410*/
411__CUDA_HOSTDEVICE_FP16_DECL__ unsigned short int __half2ushort_rz(const __half h);
412/**
413* \ingroup CUDA_MATH__HALF_MISC
414* \brief Convert a half to a signed integer in round-towards-zero mode.
415*
416* \details Convert the half-precision floating-point value \p h to a signed integer in
417* round-towards-zero mode. NaN inputs are converted to 0.
418* \param[in] h - half. Is only being read.
419*
420* \returns int
421* - \p h converted to a signed integer.
422* \internal
423* \exception-guarantee no-throw guarantee
424* \behavior reentrant, thread safe
425* \endinternal
426*/
427__CUDA_HOSTDEVICE_FP16_DECL__ int __half2int_rz(const __half h);
428/**
429* \ingroup CUDA_MATH__HALF_MISC
430* \brief Convert a half to an unsigned integer in round-towards-zero mode.
431*
432* \details Convert the half-precision floating-point value \p h to an unsigned integer
433* in round-towards-zero mode. NaN inputs are converted to 0.
434* \param[in] h - half. Is only being read.
435*
436* \returns unsigned int
437* - \p h converted to an unsigned integer.
438* \internal
439* \exception-guarantee no-throw guarantee
440* \behavior reentrant, thread safe
441* \endinternal
442*/
443__CUDA_HOSTDEVICE_FP16_DECL__ unsigned int __half2uint_rz(const __half h);
444/**
445* \ingroup CUDA_MATH__HALF_MISC
446* \brief Convert a half to a signed 64-bit integer in round-towards-zero mode.
447*
448* \details Convert the half-precision floating-point value \p h to a signed 64-bit
449* integer in round-towards-zero mode. NaN inputs return a long long int with hex value of 0x8000000000000000.
450* \param[in] h - half. Is only being read.
451*
452* \returns long long int
453* - \p h converted to a signed 64-bit integer.
454* \internal
455* \exception-guarantee no-throw guarantee
456* \behavior reentrant, thread safe
457* \endinternal
458*/
459__CUDA_HOSTDEVICE_FP16_DECL__ long long int __half2ll_rz(const __half h);
460/**
461* \ingroup CUDA_MATH__HALF_MISC
462* \brief Convert a half to an unsigned 64-bit integer in round-towards-zero
463* mode.
464*
465* \details Convert the half-precision floating-point value \p h to an unsigned 64-bit
466* integer in round-towards-zero mode. NaN inputs return 0x8000000000000000.
467* \param[in] h - half. Is only being read.
468*
469* \returns unsigned long long int
470* - \p h converted to an unsigned 64-bit integer.
471* \internal
472* \exception-guarantee no-throw guarantee
473* \behavior reentrant, thread safe
474* \endinternal
475*/
476__CUDA_HOSTDEVICE_FP16_DECL__ unsigned long long int __half2ull_rz(const __half h);
477/**
478* \ingroup CUDA_MATH__HALF_MISC
479* \brief Vector function, combines two \p __half numbers into one \p __half2 number.
480*
481* \details Combines two input \p __half number \p x and \p y into one \p __half2 number.
482* Input \p x is stored in low 16 bits of the return value, input \p y is stored
483* in high 16 bits of the return value.
484* \param[in] x - half. Is only being read.
485* \param[in] y - half. Is only being read.
486*
487* \returns __half2
488* - The \p __half2 vector with one half equal to \p x and the other to \p y.
489* \internal
490* \exception-guarantee no-throw guarantee
491* \behavior reentrant, thread safe
492* \endinternal
493*/
494__CUDA_HOSTDEVICE_FP16_DECL__ __half2 make_half2(const __half x, const __half y);
495/**
496* \ingroup CUDA_MATH__HALF_MISC
497* \brief Converts both components of float2 number to half precision in
498* round-to-nearest-even mode and returns \p half2 with converted values.
499*
500* \details Converts both components of float2 to half precision in round-to-nearest-even
501* mode and combines the results into one \p half2 number. Low 16 bits of the
502* return value correspond to \p a.x and high 16 bits of the return value
503* correspond to \p a.y.
504* \param[in] a - float2. Is only being read.
505*
506* \returns half2
507* - The \p half2 which has corresponding halves equal to the
508* converted float2 components.
509* \internal
510* \exception-guarantee no-throw guarantee
511* \behavior reentrant, thread safe
512* \endinternal
513*/
514__CUDA_HOSTDEVICE_FP16_DECL__ __half2 __float22half2_rn(const float2 a);
515/**
516* \ingroup CUDA_MATH__HALF_MISC
517* \brief Converts both halves of \p half2 to float2 and returns the result.
518*
519* \details Converts both halves of \p half2 input \p a to float2 and returns the
520* result.
521* \param[in] a - half2. Is only being read.
522*
523* \returns float2
524* - \p a converted to float2.
525* \internal
526* \exception-guarantee no-throw guarantee
527* \behavior reentrant, thread safe
528* \endinternal
529*/
530__CUDA_HOSTDEVICE_FP16_DECL__ float2 __half22float2(const __half2 a);
531#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
532/**
533* \ingroup CUDA_MATH__HALF_MISC
534* \brief Convert a half to a signed integer in round-to-nearest-even mode.
535*
536* \details Convert the half-precision floating-point value \p h to a signed integer in
537* round-to-nearest-even mode. NaN inputs are converted to 0.
538* \param[in] h - half. Is only being read.
539*
540* \returns int
541* - \p h converted to a signed integer.
542* \internal
543* \exception-guarantee no-throw guarantee
544* \behavior reentrant, thread safe
545* \endinternal
546*/
547__CUDA_FP16_DECL__ int __half2int_rn(const __half h);
548/**
549* \ingroup CUDA_MATH__HALF_MISC
550* \brief Convert a half to a signed integer in round-down mode.
551*
552* \details Convert the half-precision floating-point value \p h to a signed integer in
553* round-down mode. NaN inputs are converted to 0.
554* \param[in] h - half. Is only being read.
555*
556* \returns int
557* - \p h converted to a signed integer.
558* \internal
559* \exception-guarantee no-throw guarantee
560* \behavior reentrant, thread safe
561* \endinternal
562*/
563__CUDA_FP16_DECL__ int __half2int_rd(const __half h);
564/**
565* \ingroup CUDA_MATH__HALF_MISC
566* \brief Convert a half to a signed integer in round-up mode.
567*
568* \details Convert the half-precision floating-point value \p h to a signed integer in
569* round-up mode. NaN inputs are converted to 0.
570* \param[in] h - half. Is only being read.
571*
572* \returns int
573* - \p h converted to a signed integer.
574* \internal
575* \exception-guarantee no-throw guarantee
576* \behavior reentrant, thread safe
577* \endinternal
578*/
579__CUDA_FP16_DECL__ int __half2int_ru(const __half h);
580#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
581/**
582* \ingroup CUDA_MATH__HALF_MISC
583* \brief Convert a signed integer to a half in round-to-nearest-even mode.
584*
585* \details Convert the signed integer value \p i to a half-precision floating-point
586* value in round-to-nearest-even mode.
587* \param[in] i - int. Is only being read.
588*
589* \returns half
590* - \p i converted to half.
591* \internal
592* \exception-guarantee no-throw guarantee
593* \behavior reentrant, thread safe
594* \endinternal
595*/
596__CUDA_HOSTDEVICE_FP16_DECL__ __half __int2half_rn(const int i);
597/**
598* \ingroup CUDA_MATH__HALF_MISC
599* \brief Convert a signed integer to a half in round-towards-zero mode.
600*
601* \details Convert the signed integer value \p i to a half-precision floating-point
602* value in round-towards-zero mode.
603* \param[in] i - int. Is only being read.
604*
605* \returns half
606* - \p i converted to half.
607* \internal
608* \exception-guarantee no-throw guarantee
609* \behavior reentrant, thread safe
610* \endinternal
611*/
612__CUDA_HOSTDEVICE_FP16_DECL__ __half __int2half_rz(const int i);
613/**
614* \ingroup CUDA_MATH__HALF_MISC
615* \brief Convert a signed integer to a half in round-down mode.
616*
617* \details Convert the signed integer value \p i to a half-precision floating-point
618* value in round-down mode.
619* \param[in] i - int. Is only being read.
620*
621* \returns half
622* - \p i converted to half.
623* \internal
624* \exception-guarantee no-throw guarantee
625* \behavior reentrant, thread safe
626* \endinternal
627*/
628__CUDA_HOSTDEVICE_FP16_DECL__ __half __int2half_rd(const int i);
629/**
630* \ingroup CUDA_MATH__HALF_MISC
631* \brief Convert a signed integer to a half in round-up mode.
632*
633* \details Convert the signed integer value \p i to a half-precision floating-point
634* value in round-up mode.
635* \param[in] i - int. Is only being read.
636*
637* \returns half
638* - \p i converted to half.
639* \internal
640* \exception-guarantee no-throw guarantee
641* \behavior reentrant, thread safe
642* \endinternal
643*/
644__CUDA_HOSTDEVICE_FP16_DECL__ __half __int2half_ru(const int i);
645#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
646/**
647* \ingroup CUDA_MATH__HALF_MISC
648* \brief Convert a half to a signed short integer in round-to-nearest-even
649* mode.
650*
651* \details Convert the half-precision floating-point value \p h to a signed short
652* integer in round-to-nearest-even mode. NaN inputs are converted to 0.
653* \param[in] h - half. Is only being read.
654*
655* \returns short int
656* - \p h converted to a signed short integer.
657* \internal
658* \exception-guarantee no-throw guarantee
659* \behavior reentrant, thread safe
660* \endinternal
661*/
662__CUDA_FP16_DECL__ short int __half2short_rn(const __half h);
663/**
664* \ingroup CUDA_MATH__HALF_MISC
665* \brief Convert a half to a signed short integer in round-down mode.
666*
667* \details Convert the half-precision floating-point value \p h to a signed short
668* integer in round-down mode. NaN inputs are converted to 0.
669* \param[in] h - half. Is only being read.
670*
671* \returns short int
672* - \p h converted to a signed short integer.
673* \internal
674* \exception-guarantee no-throw guarantee
675* \behavior reentrant, thread safe
676* \endinternal
677*/
678__CUDA_FP16_DECL__ short int __half2short_rd(const __half h);
679/**
680* \ingroup CUDA_MATH__HALF_MISC
681* \brief Convert a half to a signed short integer in round-up mode.
682*
683* \details Convert the half-precision floating-point value \p h to a signed short
684* integer in round-up mode. NaN inputs are converted to 0.
685* \param[in] h - half. Is only being read.
686*
687* \returns short int
688* - \p h converted to a signed short integer.
689* \internal
690* \exception-guarantee no-throw guarantee
691* \behavior reentrant, thread safe
692* \endinternal
693*/
694__CUDA_FP16_DECL__ short int __half2short_ru(const __half h);
695#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
696/**
697* \ingroup CUDA_MATH__HALF_MISC
698* \brief Convert a signed short integer to a half in round-to-nearest-even
699* mode.
700*
701* \details Convert the signed short integer value \p i to a half-precision floating-point
702* value in round-to-nearest-even mode.
703* \param[in] i - short int. Is only being read.
704*
705* \returns half
706* - \p i converted to half.
707* \internal
708* \exception-guarantee no-throw guarantee
709* \behavior reentrant, thread safe
710* \endinternal
711*/
712__CUDA_HOSTDEVICE_FP16_DECL__ __half __short2half_rn(const short int i);
713/**
714* \ingroup CUDA_MATH__HALF_MISC
715* \brief Convert a signed short integer to a half in round-towards-zero mode.
716*
717* \details Convert the signed short integer value \p i to a half-precision floating-point
718* value in round-towards-zero mode.
719* \param[in] i - short int. Is only being read.
720*
721* \returns half
722* - \p i converted to half.
723* \internal
724* \exception-guarantee no-throw guarantee
725* \behavior reentrant, thread safe
726* \endinternal
727*/
728__CUDA_HOSTDEVICE_FP16_DECL__ __half __short2half_rz(const short int i);
729/**
730* \ingroup CUDA_MATH__HALF_MISC
731* \brief Convert a signed short integer to a half in round-down mode.
732*
733* \details Convert the signed short integer value \p i to a half-precision floating-point
734* value in round-down mode.
735* \param[in] i - short int. Is only being read.
736*
737* \returns half
738* - \p i converted to half.
739* \internal
740* \exception-guarantee no-throw guarantee
741* \behavior reentrant, thread safe
742* \endinternal
743*/
744__CUDA_HOSTDEVICE_FP16_DECL__ __half __short2half_rd(const short int i);
745/**
746* \ingroup CUDA_MATH__HALF_MISC
747* \brief Convert a signed short integer to a half in round-up mode.
748*
749* \details Convert the signed short integer value \p i to a half-precision floating-point
750* value in round-up mode.
751* \param[in] i - short int. Is only being read.
752*
753* \returns half
754* - \p i converted to half.
755* \internal
756* \exception-guarantee no-throw guarantee
757* \behavior reentrant, thread safe
758* \endinternal
759*/
760__CUDA_HOSTDEVICE_FP16_DECL__ __half __short2half_ru(const short int i);
761#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
762/**
763* \ingroup CUDA_MATH__HALF_MISC
764* \brief Convert a half to an unsigned integer in round-to-nearest-even mode.
765*
766* \details Convert the half-precision floating-point value \p h to an unsigned integer
767* in round-to-nearest-even mode. NaN inputs are converted to 0.
768* \param[in] h - half. Is only being read.
769*
770* \returns unsigned int
771* - \p h converted to an unsigned integer.
772* \internal
773* \exception-guarantee no-throw guarantee
774* \behavior reentrant, thread safe
775* \endinternal
776*/
777__CUDA_FP16_DECL__ unsigned int __half2uint_rn(const __half h);
778/**
779* \ingroup CUDA_MATH__HALF_MISC
780* \brief Convert a half to an unsigned integer in round-down mode.
781*
782* \details Convert the half-precision floating-point value \p h to an unsigned integer
783* in round-down mode. NaN inputs are converted to 0.
784* \param[in] h - half. Is only being read.
785*
786* \returns unsigned int
787* - \p h converted to an unsigned integer.
788* \internal
789* \exception-guarantee no-throw guarantee
790* \behavior reentrant, thread safe
791* \endinternal
792*/
793__CUDA_FP16_DECL__ unsigned int __half2uint_rd(const __half h);
794/**
795* \ingroup CUDA_MATH__HALF_MISC
796* \brief Convert a half to an unsigned integer in round-up mode.
797*
798* \details Convert the half-precision floating-point value \p h to an unsigned integer
799* in round-up mode. NaN inputs are converted to 0.
800* \param[in] h - half. Is only being read.
801*
802* \returns unsigned int
803* - \p h converted to an unsigned integer.
804* \internal
805* \exception-guarantee no-throw guarantee
806* \behavior reentrant, thread safe
807* \endinternal
808*/
809__CUDA_FP16_DECL__ unsigned int __half2uint_ru(const __half h);
810#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
811/**
812* \ingroup CUDA_MATH__HALF_MISC
813* \brief Convert an unsigned integer to a half in round-to-nearest-even mode.
814*
815* \details Convert the unsigned integer value \p i to a half-precision floating-point
816* value in round-to-nearest-even mode.
817* \param[in] i - unsigned int. Is only being read.
818*
819* \returns half
820* - \p i converted to half.
821* \internal
822* \exception-guarantee no-throw guarantee
823* \behavior reentrant, thread safe
824* \endinternal
825*/
826__CUDA_HOSTDEVICE_FP16_DECL__ __half __uint2half_rn(const unsigned int i);
827/**
828* \ingroup CUDA_MATH__HALF_MISC
829* \brief Convert an unsigned integer to a half in round-towards-zero mode.
830*
831* \details Convert the unsigned integer value \p i to a half-precision floating-point
832* value in round-towards-zero mode.
833* \param[in] i - unsigned int. Is only being read.
834*
835* \returns half
836* - \p i converted to half.
837* \internal
838* \exception-guarantee no-throw guarantee
839* \behavior reentrant, thread safe
840* \endinternal
841*/
842__CUDA_HOSTDEVICE_FP16_DECL__ __half __uint2half_rz(const unsigned int i);
843/**
844* \ingroup CUDA_MATH__HALF_MISC
845* \brief Convert an unsigned integer to a half in round-down mode.
846*
847* \details Convert the unsigned integer value \p i to a half-precision floating-point
848* value in round-down mode.
849* \param[in] i - unsigned int. Is only being read.
850*
851* \returns half
852* - \p i converted to half.
853* \internal
854* \exception-guarantee no-throw guarantee
855* \behavior reentrant, thread safe
856* \endinternal
857*/
858__CUDA_HOSTDEVICE_FP16_DECL__ __half __uint2half_rd(const unsigned int i);
859/**
860* \ingroup CUDA_MATH__HALF_MISC
861* \brief Convert an unsigned integer to a half in round-up mode.
862*
863* \details Convert the unsigned integer value \p i to a half-precision floating-point
864* value in round-up mode.
865* \param[in] i - unsigned int. Is only being read.
866*
867* \returns half
868* - \p i converted to half.
869* \internal
870* \exception-guarantee no-throw guarantee
871* \behavior reentrant, thread safe
872* \endinternal
873*/
874__CUDA_HOSTDEVICE_FP16_DECL__ __half __uint2half_ru(const unsigned int i);
875#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
876/**
877* \ingroup CUDA_MATH__HALF_MISC
878* \brief Convert a half to an unsigned short integer in round-to-nearest-even
879* mode.
880*
881* \details Convert the half-precision floating-point value \p h to an unsigned short
882* integer in round-to-nearest-even mode. NaN inputs are converted to 0.
883* \param[in] h - half. Is only being read.
884*
885* \returns unsigned short int
886* - \p h converted to an unsigned short integer.
887* \internal
888* \exception-guarantee no-throw guarantee
889* \behavior reentrant, thread safe
890* \endinternal
891*/
892__CUDA_FP16_DECL__ unsigned short int __half2ushort_rn(const __half h);
893/**
894* \ingroup CUDA_MATH__HALF_MISC
895* \brief Convert a half to an unsigned short integer in round-down mode.
896*
897* \details Convert the half-precision floating-point value \p h to an unsigned short
898* integer in round-down mode. NaN inputs are converted to 0.
899* \param[in] h - half. Is only being read.
900*
901* \returns unsigned short int
902* - \p h converted to an unsigned short integer.
903*/
904__CUDA_FP16_DECL__ unsigned short int __half2ushort_rd(const __half h);
905/**
906* \ingroup CUDA_MATH__HALF_MISC
907* \brief Convert a half to an unsigned short integer in round-up mode.
908*
909* \details Convert the half-precision floating-point value \p h to an unsigned short
910* integer in round-up mode. NaN inputs are converted to 0.
911* \param[in] h - half. Is only being read.
912*
913* \returns unsigned short int
914* - \p h converted to an unsigned short integer.
915*/
916__CUDA_FP16_DECL__ unsigned short int __half2ushort_ru(const __half h);
917#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
918/**
919* \ingroup CUDA_MATH__HALF_MISC
920* \brief Convert an unsigned short integer to a half in round-to-nearest-even
921* mode.
922*
923* \details Convert the unsigned short integer value \p i to a half-precision floating-point
924* value in round-to-nearest-even mode.
925* \param[in] i - unsigned short int. Is only being read.
926*
927* \returns half
928* - \p i converted to half.
929* \internal
930* \exception-guarantee no-throw guarantee
931* \behavior reentrant, thread safe
932* \endinternal
933*/
934__CUDA_HOSTDEVICE_FP16_DECL__ __half __ushort2half_rn(const unsigned short int i);
935/**
936* \ingroup CUDA_MATH__HALF_MISC
937* \brief Convert an unsigned short integer to a half in round-towards-zero
938* mode.
939*
940* \details Convert the unsigned short integer value \p i to a half-precision floating-point
941* value in round-towards-zero mode.
942* \param[in] i - unsigned short int. Is only being read.
943*
944* \returns half
945* - \p i converted to half.
946* \internal
947* \exception-guarantee no-throw guarantee
948* \behavior reentrant, thread safe
949* \endinternal
950*/
951__CUDA_HOSTDEVICE_FP16_DECL__ __half __ushort2half_rz(const unsigned short int i);
952/**
953* \ingroup CUDA_MATH__HALF_MISC
954* \brief Convert an unsigned short integer to a half in round-down mode.
955*
956* \details Convert the unsigned short integer value \p i to a half-precision floating-point
957* value in round-down mode.
958* \param[in] i - unsigned short int. Is only being read.
959*
960* \returns half
961* - \p i converted to half.
962* \internal
963* \exception-guarantee no-throw guarantee
964* \behavior reentrant, thread safe
965* \endinternal
966*/
967__CUDA_HOSTDEVICE_FP16_DECL__ __half __ushort2half_rd(const unsigned short int i);
968/**
969* \ingroup CUDA_MATH__HALF_MISC
970* \brief Convert an unsigned short integer to a half in round-up mode.
971*
972* \details Convert the unsigned short integer value \p i to a half-precision floating-point
973* value in round-up mode.
974* \param[in] i - unsigned short int. Is only being read.
975*
976* \returns half
977* - \p i converted to half.
978* \internal
979* \exception-guarantee no-throw guarantee
980* \behavior reentrant, thread safe
981* \endinternal
982*/
983__CUDA_HOSTDEVICE_FP16_DECL__ __half __ushort2half_ru(const unsigned short int i);
984#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
985/**
986* \ingroup CUDA_MATH__HALF_MISC
987* \brief Convert a half to an unsigned 64-bit integer in round-to-nearest-even
988* mode.
989*
990* \details Convert the half-precision floating-point value \p h to an unsigned 64-bit
991* integer in round-to-nearest-even mode. NaN inputs return 0x8000000000000000.
992* \param[in] h - half. Is only being read.
993*
994* \returns unsigned long long int
995* - \p h converted to an unsigned 64-bit integer.
996* \internal
997* \exception-guarantee no-throw guarantee
998* \behavior reentrant, thread safe
999* \endinternal
1000*/
1001__CUDA_FP16_DECL__ unsigned long long int __half2ull_rn(const __half h);
1002/**
1003* \ingroup CUDA_MATH__HALF_MISC
1004* \brief Convert a half to an unsigned 64-bit integer in round-down mode.
1005*
1006* \details Convert the half-precision floating-point value \p h to an unsigned 64-bit
1007* integer in round-down mode. NaN inputs return 0x8000000000000000.
1008* \param[in] h - half. Is only being read.
1009*
1010* \returns unsigned long long int
1011* - \p h converted to an unsigned 64-bit integer.
1012* \internal
1013* \exception-guarantee no-throw guarantee
1014* \behavior reentrant, thread safe
1015* \endinternal
1016*/
1017__CUDA_FP16_DECL__ unsigned long long int __half2ull_rd(const __half h);
1018/**
1019* \ingroup CUDA_MATH__HALF_MISC
1020* \brief Convert a half to an unsigned 64-bit integer in round-up mode.
1021*
1022* \details Convert the half-precision floating-point value \p h to an unsigned 64-bit
1023* integer in round-up mode. NaN inputs return 0x8000000000000000.
1024* \param[in] h - half. Is only being read.
1025*
1026* \returns unsigned long long int
1027* - \p h converted to an unsigned 64-bit integer.
1028* \internal
1029* \exception-guarantee no-throw guarantee
1030* \behavior reentrant, thread safe
1031* \endinternal
1032*/
1033__CUDA_FP16_DECL__ unsigned long long int __half2ull_ru(const __half h);
1034#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
1035/**
1036* \ingroup CUDA_MATH__HALF_MISC
1037* \brief Convert an unsigned 64-bit integer to a half in round-to-nearest-even
1038* mode.
1039*
1040* \details Convert the unsigned 64-bit integer value \p i to a half-precision floating-point
1041* value in round-to-nearest-even mode.
1042* \param[in] i - unsigned long long int. Is only being read.
1043*
1044* \returns half
1045* - \p i converted to half.
1046* \internal
1047* \exception-guarantee no-throw guarantee
1048* \behavior reentrant, thread safe
1049* \endinternal
1050*/
1051__CUDA_HOSTDEVICE_FP16_DECL__ __half __ull2half_rn(const unsigned long long int i);
1052/**
1053* \ingroup CUDA_MATH__HALF_MISC
1054* \brief Convert an unsigned 64-bit integer to a half in round-towards-zero
1055* mode.
1056*
1057* \details Convert the unsigned 64-bit integer value \p i to a half-precision floating-point
1058* value in round-towards-zero mode.
1059* \param[in] i - unsigned long long int. Is only being read.
1060*
1061* \returns half
1062* - \p i converted to half.
1063* \internal
1064* \exception-guarantee no-throw guarantee
1065* \behavior reentrant, thread safe
1066* \endinternal
1067*/
1068__CUDA_HOSTDEVICE_FP16_DECL__ __half __ull2half_rz(const unsigned long long int i);
1069/**
1070* \ingroup CUDA_MATH__HALF_MISC
1071* \brief Convert an unsigned 64-bit integer to a half in round-down mode.
1072*
1073* \details Convert the unsigned 64-bit integer value \p i to a half-precision floating-point
1074* value in round-down mode.
1075* \param[in] i - unsigned long long int. Is only being read.
1076*
1077* \returns half
1078* - \p i converted to half.
1079* \internal
1080* \exception-guarantee no-throw guarantee
1081* \behavior reentrant, thread safe
1082* \endinternal
1083*/
1084__CUDA_HOSTDEVICE_FP16_DECL__ __half __ull2half_rd(const unsigned long long int i);
1085/**
1086* \ingroup CUDA_MATH__HALF_MISC
1087* \brief Convert an unsigned 64-bit integer to a half in round-up mode.
1088*
1089* \details Convert the unsigned 64-bit integer value \p i to a half-precision floating-point
1090* value in round-up mode.
1091* \param[in] i - unsigned long long int. Is only being read.
1092*
1093* \returns half
1094* - \p i converted to half.
1095* \internal
1096* \exception-guarantee no-throw guarantee
1097* \behavior reentrant, thread safe
1098* \endinternal
1099*/
1100__CUDA_HOSTDEVICE_FP16_DECL__ __half __ull2half_ru(const unsigned long long int i);
1101#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
1102/**
1103* \ingroup CUDA_MATH__HALF_MISC
1104* \brief Convert a half to a signed 64-bit integer in round-to-nearest-even
1105* mode.
1106*
1107* \details Convert the half-precision floating-point value \p h to a signed 64-bit
1108* integer in round-to-nearest-even mode. NaN inputs return a long long int with hex value of 0x8000000000000000.
1109* \param[in] h - half. Is only being read.
1110*
1111* \returns long long int
1112* - \p h converted to a signed 64-bit integer.
1113* \internal
1114* \exception-guarantee no-throw guarantee
1115* \behavior reentrant, thread safe
1116* \endinternal
1117*/
1118__CUDA_FP16_DECL__ long long int __half2ll_rn(const __half h);
1119/**
1120* \ingroup CUDA_MATH__HALF_MISC
1121* \brief Convert a half to a signed 64-bit integer in round-down mode.
1122*
1123* \details Convert the half-precision floating-point value \p h to a signed 64-bit
1124* integer in round-down mode. NaN inputs return a long long int with hex value of 0x8000000000000000.
1125* \param[in] h - half. Is only being read.
1126*
1127* \returns long long int
1128* - \p h converted to a signed 64-bit integer.
1129* \internal
1130* \exception-guarantee no-throw guarantee
1131* \behavior reentrant, thread safe
1132* \endinternal
1133*/
1134__CUDA_FP16_DECL__ long long int __half2ll_rd(const __half h);
1135/**
1136* \ingroup CUDA_MATH__HALF_MISC
1137* \brief Convert a half to a signed 64-bit integer in round-up mode.
1138*
1139* \details Convert the half-precision floating-point value \p h to a signed 64-bit
1140* integer in round-up mode. NaN inputs return a long long int with hex value of 0x8000000000000000.
1141* \param[in] h - half. Is only being read.
1142*
1143* \returns long long int
1144* - \p h converted to a signed 64-bit integer.
1145* \internal
1146* \exception-guarantee no-throw guarantee
1147* \behavior reentrant, thread safe
1148* \endinternal
1149*/
1150__CUDA_FP16_DECL__ long long int __half2ll_ru(const __half h);
1151#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
1152/**
1153* \ingroup CUDA_MATH__HALF_MISC
1154* \brief Convert a signed 64-bit integer to a half in round-to-nearest-even
1155* mode.
1156*
1157* \details Convert the signed 64-bit integer value \p i to a half-precision floating-point
1158* value in round-to-nearest-even mode.
1159* \param[in] i - long long int. Is only being read.
1160*
1161* \returns half
1162* - \p i converted to half.
1163* \internal
1164* \exception-guarantee no-throw guarantee
1165* \behavior reentrant, thread safe
1166* \endinternal
1167*/
1168__CUDA_HOSTDEVICE_FP16_DECL__ __half __ll2half_rn(const long long int i);
1169/**
1170* \ingroup CUDA_MATH__HALF_MISC
1171* \brief Convert a signed 64-bit integer to a half in round-towards-zero mode.
1172*
1173* \details Convert the signed 64-bit integer value \p i to a half-precision floating-point
1174* value in round-towards-zero mode.
1175* \param[in] i - long long int. Is only being read.
1176*
1177* \returns half
1178* - \p i converted to half.
1179*/
1180__CUDA_HOSTDEVICE_FP16_DECL__ __half __ll2half_rz(const long long int i);
1181/**
1182* \ingroup CUDA_MATH__HALF_MISC
1183* \brief Convert a signed 64-bit integer to a half in round-down mode.
1184*
1185* \details Convert the signed 64-bit integer value \p i to a half-precision floating-point
1186* value in round-down mode.
1187* \param[in] i - long long int. Is only being read.
1188*
1189* \returns half
1190* - \p i converted to half.
1191* \internal
1192* \exception-guarantee no-throw guarantee
1193* \behavior reentrant, thread safe
1194* \endinternal
1195*/
1196__CUDA_HOSTDEVICE_FP16_DECL__ __half __ll2half_rd(const long long int i);
1197/**
1198* \ingroup CUDA_MATH__HALF_MISC
1199* \brief Convert a signed 64-bit integer to a half in round-up mode.
1200*