codekingpro/portable-devtools
114k
1/*
2* Copyright 1993-2023 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#if !defined(__CUDA_FP16_HPP__)
51#define __CUDA_FP16_HPP__
52
53#if !defined(__CUDA_FP16_H__)
54#error "Do not include this file directly. Instead, include cuda_fp16.h."
55#endif
56
57#if !defined(IF_DEVICE_OR_CUDACC)
58#if defined(__CUDACC__)
59 #define IF_DEVICE_OR_CUDACC(d, c, f) NV_IF_ELSE_TARGET(NV_IS_DEVICE, d, c)
60#else
61 #define IF_DEVICE_OR_CUDACC(d, c, f) NV_IF_ELSE_TARGET(NV_IS_DEVICE, d, f)
62#endif
63#endif
64
65/* Macros for half & half2 binary arithmetic */
66#define __BINARY_OP_HALF_MACRO(name) /* do */ {\
67 __half val; \
68 asm( "{" __CUDA_FP16_STRINGIFY(name) ".f16 %0,%1,%2;\n}" \
69 :"=h"(__HALF_TO_US(val)) : "h"(__HALF_TO_CUS(a)),"h"(__HALF_TO_CUS(b))); \
70 return val; \
71} /* while(0) */
72#define __BINARY_OP_HALF2_MACRO(name) /* do */ {\
73 __half2 val; \
74 asm( "{" __CUDA_FP16_STRINGIFY(name) ".f16x2 %0,%1,%2;\n}" \
75 :"=r"(__HALF2_TO_UI(val)) : "r"(__HALF2_TO_CUI(a)),"r"(__HALF2_TO_CUI(b))); \
76 return val; \
77} /* while(0) */
78#define __TERNARY_OP_HALF_MACRO(name) /* do */ {\
79 __half val; \
80 asm( "{" __CUDA_FP16_STRINGIFY(name) ".f16 %0,%1,%2,%3;\n}" \
81 :"=h"(__HALF_TO_US(val)) : "h"(__HALF_TO_CUS(a)),"h"(__HALF_TO_CUS(b)),"h"(__HALF_TO_CUS(c))); \
82 return val; \
83} /* while(0) */
84#define __TERNARY_OP_HALF2_MACRO(name) /* do */ {\
85 __half2 val; \
86 asm( "{" __CUDA_FP16_STRINGIFY(name) ".f16x2 %0,%1,%2,%3;\n}" \
87 :"=r"(__HALF2_TO_UI(val)) : "r"(__HALF2_TO_CUI(a)),"r"(__HALF2_TO_CUI(b)),"r"(__HALF2_TO_CUI(c))); \
88 return val; \
89} /* while(0) */
90
91/* All other definitions in this file are only visible to C++ compilers */
92#if defined(__cplusplus)
93
94/**
95 * \ingroup CUDA_MATH_INTRINSIC_HALF_CONSTANTS
96 * \brief Defines floating-point positive infinity value for the \p half data type
97 */
98#define CUDART_INF_FP16 __ushort_as_half((unsigned short)0x7C00U)
99/**
100 * \ingroup CUDA_MATH_INTRINSIC_HALF_CONSTANTS
101 * \brief Defines canonical NaN value for the \p half data type
102 */
103#define CUDART_NAN_FP16 __ushort_as_half((unsigned short)0x7FFFU)
104/**
105 * \ingroup CUDA_MATH_INTRINSIC_HALF_CONSTANTS
106 * \brief Defines a minimum representable (denormalized) value for the \p half data type
107 */
108#define CUDART_MIN_DENORM_FP16 __ushort_as_half((unsigned short)0x0001U)
109/**
110 * \ingroup CUDA_MATH_INTRINSIC_HALF_CONSTANTS
111 * \brief Defines a maximum representable value for the \p half data type
112 */
113#define CUDART_MAX_NORMAL_FP16 __ushort_as_half((unsigned short)0x7BFFU)
114/**
115 * \ingroup CUDA_MATH_INTRINSIC_HALF_CONSTANTS
116 * \brief Defines a negative zero value for the \p half data type
117 */
118#define CUDART_NEG_ZERO_FP16 __ushort_as_half((unsigned short)0x8000U)
119/**
120 * \ingroup CUDA_MATH_INTRINSIC_HALF_CONSTANTS
121 * \brief Defines a positive zero value for the \p half data type
122 */
123#define CUDART_ZERO_FP16 __ushort_as_half((unsigned short)0x0000U)
124/**
125 * \ingroup CUDA_MATH_INTRINSIC_HALF_CONSTANTS
126 * \brief Defines a value of 1.0 for the \p half data type
127 */
128#define CUDART_ONE_FP16 __ushort_as_half((unsigned short)0x3C00U)
129
130__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half &__half::operator=(const __half_raw &hr) { __x = hr.x; return *this; }
131__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ volatile __half &__half::operator=(const __half_raw &hr) volatile { __x = hr.x; return *this; }
132__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ volatile __half &__half::operator=(const volatile __half_raw &hr) volatile { __x = hr.x; return *this; }
133__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator __half_raw() const { __half_raw ret; ret.x = __x; return ret; }
134__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator __half_raw() const volatile { __half_raw ret; ret.x = __x; return ret; }
135#if !defined(__CUDA_NO_HALF_CONVERSIONS__)
136__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator float() const { return __half2float(*this); }
137__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half &__half::operator=(const float f) { __x = __float2half(f).__x; return *this; }
138__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half &__half::operator=(const double f) { __x = __double2half(f).__x; return *this; }
139#if !(defined __CUDA_FP16_DISABLE_IMPLICIT_INTEGER_CONVERTS_FOR_HOST_COMPILERS__) || (defined __CUDACC__)
140__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator signed char() const { return __half2char_rz(*this); }
141__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator unsigned char() const { return __half2uchar_rz(*this); }
142__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator char() const {
143 char value;
144 /* Suppress VS warning: warning C4127: conditional expression is constant */
145#if defined(_MSC_VER) && !defined(__CUDA_ARCH__)
146#pragma warning (push)
147#pragma warning (disable: 4127)
148#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */
149 if (((char)-1) < (char)0)
150#if defined(_MSC_VER) && !defined(__CUDA_ARCH__)
151#pragma warning (pop)
152#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */
153 {
154 value = static_cast<char>(__half2char_rz(*this));
155 }
156 else
157 {
158 value = static_cast<char>(__half2uchar_rz(*this));
159 }
160 return value;
161}
162__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator short() const { return __half2short_rz(*this); }
163__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator unsigned short() const { return __half2ushort_rz(*this); }
164__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator int() const { return __half2int_rz(*this); }
165__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator unsigned int() const { return __half2uint_rz(*this); }
166__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator long() const {
167 long retval;
168 /* Suppress VS warning: warning C4127: conditional expression is constant */
169#if defined(_MSC_VER) && !defined(__CUDA_ARCH__)
170#pragma warning (push)
171#pragma warning (disable: 4127)
172#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */
173 if (sizeof(long) == sizeof(long long))
174#if defined(_MSC_VER) && !defined(__CUDA_ARCH__)
175#pragma warning (pop)
176#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */
177 {
178 retval = static_cast<long>(__half2ll_rz(*this));
179 }
180 else
181 {
182 retval = static_cast<long>(__half2int_rz(*this));
183 }
184 return retval;
185}
186__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator unsigned long() const {
187 unsigned long retval;
188 /* Suppress VS warning: warning C4127: conditional expression is constant */
189#if defined(_MSC_VER) && !defined(__CUDA_ARCH__)
190#pragma warning (push)
191#pragma warning (disable: 4127)
192#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */
193 if (sizeof(unsigned long) == sizeof(unsigned long long))
194#if defined(_MSC_VER) && !defined(__CUDA_ARCH__)
195#pragma warning (pop)
196#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */
197 {
198 retval = static_cast<unsigned long>(__half2ull_rz(*this));
199 }
200 else
201 {
202 retval = static_cast<unsigned long>(__half2uint_rz(*this));
203 }
204 return retval;
205}
206__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator long long() const { return __half2ll_rz(*this); }
207__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half::operator unsigned long long() const { return __half2ull_rz(*this); }
208__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half &__half::operator=(const short val) { __x = __short2half_rn(val).__x; return *this; }
209__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half &__half::operator=(const unsigned short val) { __x = __ushort2half_rn(val).__x; return *this; }
210__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half &__half::operator=(const int val) { __x = __int2half_rn(val).__x; return *this; }
211__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half &__half::operator=(const unsigned int val) { __x = __uint2half_rn(val).__x; return *this; }
212__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half &__half::operator=(const long long val) { __x = __ll2half_rn(val).__x; return *this; }
213__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half &__half::operator=(const unsigned long long val) { __x = __ull2half_rn(val).__x; return *this; }
214#endif /* #if !(defined __CUDA_FP16_DISABLE_IMPLICIT_INTEGER_CONVERTS_FOR_HOST_COMPILERS__) || (defined __CUDACC__) */
215#endif /* !defined(__CUDA_NO_HALF_CONVERSIONS__) */
216#if !defined(__CUDA_NO_HALF_OPERATORS__)
217/* Some basic arithmetic operations expected of a builtin */
218__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half operator+(const __half &lh, const __half &rh) { return __hadd(lh, rh); }
219__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half operator-(const __half &lh, const __half &rh) { return __hsub(lh, rh); }
220__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half operator*(const __half &lh, const __half &rh) { return __hmul(lh, rh); }
221__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half operator/(const __half &lh, const __half &rh) { return __hdiv(lh, rh); }
222__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half &operator+=(__half &lh, const __half &rh) { lh = __hadd(lh, rh); return lh; }
223__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half &operator-=(__half &lh, const __half &rh) { lh = __hsub(lh, rh); return lh; }
224__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half &operator*=(__half &lh, const __half &rh) { lh = __hmul(lh, rh); return lh; }
225__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half &operator/=(__half &lh, const __half &rh) { lh = __hdiv(lh, rh); return lh; }
226__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half &operator++(__half &h) { __half_raw one; one.x = 0x3C00U; h += one; return h; }
227__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half &operator--(__half &h) { __half_raw one; one.x = 0x3C00U; h -= one; return h; }
228__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half operator++(__half &h, const int ignored)
229{
230 // ignored on purpose. Parameter only needed to distinguish the function declaration from other types of operators.
231 static_cast<void>(ignored);
232
233 const __half ret = h;
234 __half_raw one;
235 one.x = 0x3C00U;
236 h += one;
237 return ret;
238}
239__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half operator--(__half &h, const int ignored)
240{
241 // ignored on purpose. Parameter only needed to distinguish the function declaration from other types of operators.
242 static_cast<void>(ignored);
243
244 const __half ret = h;
245 __half_raw one;
246 one.x = 0x3C00U;
247 h -= one;
248 return ret;
249}
250__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half operator+(const __half &h) { return h; }
251__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half operator-(const __half &h) { return __hneg(h); }
252__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator==(const __half &lh, const __half &rh) { return __heq(lh, rh); }
253__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator!=(const __half &lh, const __half &rh) { return __hneu(lh, rh); }
254__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator> (const __half &lh, const __half &rh) { return __hgt(lh, rh); }
255__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator< (const __half &lh, const __half &rh) { return __hlt(lh, rh); }
256__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator>=(const __half &lh, const __half &rh) { return __hge(lh, rh); }
257__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator<=(const __half &lh, const __half &rh) { return __hle(lh, rh); }
258#endif /* !defined(__CUDA_NO_HALF_OPERATORS__) */
259#if defined(__CPP_VERSION_AT_LEAST_11_FP16)
260__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half2 &__half2::operator=(const __half2 &&src) {
261NV_IF_ELSE_TARGET(NV_IS_DEVICE,
262 __HALF2_TO_UI(*this) = std::move(__HALF2_TO_CUI(src));
263,
264 this->x = src.x;
265 this->y = src.y;
266)
267 return *this;
268}
269#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP16) */
270__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half2 &__half2::operator=(const __half2 &src) {
271NV_IF_ELSE_TARGET(NV_IS_DEVICE,
272 __HALF2_TO_UI(*this) = __HALF2_TO_CUI(src);
273,
274 this->x = src.x;
275 this->y = src.y;
276)
277 return *this;
278}
279__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half2 &__half2::operator=(const __half2_raw &h2r) {
280NV_IF_ELSE_TARGET(NV_IS_DEVICE,
281 __HALF2_TO_UI(*this) = __HALF2_TO_CUI(h2r);
282,
283 __half_raw tr;
284 tr.x = h2r.x;
285 this->x = static_cast<__half>(tr);
286 tr.x = h2r.y;
287 this->y = static_cast<__half>(tr);
288)
289 return *this;
290}
291__CUDA_HOSTDEVICE__ __CUDA_FP16_INLINE__ __half2::operator __half2_raw() const {
292 __half2_raw ret;
293NV_IF_ELSE_TARGET(NV_IS_DEVICE,
294 ret.x = 0U;
295 ret.y = 0U;
296 __HALF2_TO_UI(ret) = __HALF2_TO_CUI(*this);
297,
298 ret.x = static_cast<__half_raw>(this->x).x;
299 ret.y = static_cast<__half_raw>(this->y).x;
300)
301 return ret;
302}
303#if !defined(__CUDA_NO_HALF2_OPERATORS__)
304__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2 operator+(const __half2 &lh, const __half2 &rh) { return __hadd2(lh, rh); }
305__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2 operator-(const __half2 &lh, const __half2 &rh) { return __hsub2(lh, rh); }
306__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2 operator*(const __half2 &lh, const __half2 &rh) { return __hmul2(lh, rh); }
307__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2 operator/(const __half2 &lh, const __half2 &rh) { return __h2div(lh, rh); }
308__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2& operator+=(__half2 &lh, const __half2 &rh) { lh = __hadd2(lh, rh); return lh; }
309__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2& operator-=(__half2 &lh, const __half2 &rh) { lh = __hsub2(lh, rh); return lh; }
310__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2& operator*=(__half2 &lh, const __half2 &rh) { lh = __hmul2(lh, rh); return lh; }
311__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2& operator/=(__half2 &lh, const __half2 &rh) { lh = __h2div(lh, rh); return lh; }
312__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2 &operator++(__half2 &h) { __half2_raw one; one.x = 0x3C00U; one.y = 0x3C00U; h = __hadd2(h, one); return h; }
313__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2 &operator--(__half2 &h) { __half2_raw one; one.x = 0x3C00U; one.y = 0x3C00U; h = __hsub2(h, one); return h; }
314__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2 operator++(__half2 &h, const int ignored)
315{
316 // ignored on purpose. Parameter only needed to distinguish the function declaration from other types of operators.
317 static_cast<void>(ignored);
318
319 const __half2 ret = h;
320 __half2_raw one;
321 one.x = 0x3C00U;
322 one.y = 0x3C00U;
323 h = __hadd2(h, one);
324 return ret;
325}
326__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2 operator--(__half2 &h, const int ignored)
327{
328 // ignored on purpose. Parameter only needed to distinguish the function declaration from other types of operators.
329 static_cast<void>(ignored);
330
331 const __half2 ret = h;
332 __half2_raw one;
333 one.x = 0x3C00U;
334 one.y = 0x3C00U;
335 h = __hsub2(h, one);
336 return ret;
337}
338__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2 operator+(const __half2 &h) { return h; }
339__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ __half2 operator-(const __half2 &h) { return __hneg2(h); }
340__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator==(const __half2 &lh, const __half2 &rh) { return __hbeq2(lh, rh); }
341__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator!=(const __half2 &lh, const __half2 &rh) { return __hbneu2(lh, rh); }
342__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator>(const __half2 &lh, const __half2 &rh) { return __hbgt2(lh, rh); }
343__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator<(const __half2 &lh, const __half2 &rh) { return __hblt2(lh, rh); }
344__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator>=(const __half2 &lh, const __half2 &rh) { return __hbge2(lh, rh); }
345__CUDA_HOSTDEVICE__ __CUDA_FP16_FORCEINLINE__ bool operator<=(const __half2 &lh, const __half2 &rh) { return __hble2(lh, rh); }
346#endif /* !defined(__CUDA_NO_HALF2_OPERATORS__) */
347
348/* Restore warning for multiple assignment operators */
349#if defined(_MSC_VER) && _MSC_VER >= 1500
350#pragma warning( pop )
351#endif /* defined(_MSC_VER) && _MSC_VER >= 1500 */
352
353/* Restore -Weffc++ warnings from here on */
354#if defined(__GNUC__)
355#if __GNUC__ > 4 || (__GNUC__ == 4 && __GNUC_MINOR__ >= 6)
356#pragma GCC diagnostic pop
357#endif /* __GNUC__ > 4 || (__GNUC__ == 4 && __GNUC_MINOR__ >= 6) */
358#endif /* defined(__GNUC__) */
359
360#undef __CUDA_HOSTDEVICE__
361#undef __CUDA_ALIGN__
362
363#ifndef __CUDACC_RTC__ /* no host functions in NVRTC mode */
364static inline unsigned short __internal_float2half(const float f, unsigned int &sign, unsigned int &remainder)
365{
366 unsigned int x;
367 unsigned int u;
368 unsigned int result;
369#if defined(__CUDACC__)
370 (void)memcpy(&x, &f, sizeof(f));
371#else
372 (void)std::memcpy(&x, &f, sizeof(f));
373#endif
374 u = (x & 0x7fffffffU);
375 sign = ((x >> 16U) & 0x8000U);
376 // NaN/+Inf/-Inf
377 if (u >= 0x7f800000U) {
378 remainder = 0U;
379 result = ((u == 0x7f800000U) ? (sign | 0x7c00U) : 0x7fffU);
380 } else if (u > 0x477fefffU) { // Overflows
381 remainder = 0x80000000U;
382 result = (sign | 0x7bffU);
383 } else if (u >= 0x38800000U) { // Normal numbers
384 remainder = u << 19U;
385 u -= 0x38000000U;
386 result = (sign | (u >> 13U));
387 } else if (u < 0x33000001U) { // +0/-0
388 remainder = u;
389 result = sign;
390 } else { // Denormal numbers
391 const unsigned int exponent = u >> 23U;
392 const unsigned int shift = 0x7eU - exponent;
393 unsigned int mantissa = (u & 0x7fffffU);
394 mantissa |= 0x800000U;
395 remainder = mantissa << (32U - shift);
396 result = (sign | (mantissa >> shift));
397 result &= 0x0000FFFFU;
398 }
399 return static_cast<unsigned short>(result);
400}
401#endif /* #if !defined(__CUDACC_RTC__) */
402
403__CUDA_HOSTDEVICE_FP16_DECL__ __half __double2half(const double a)
404{
405IF_DEVICE_OR_CUDACC(
406 __half val;
407 asm("{ cvt.rn.f16.f64 %0, %1;}\n" : "=h"(__HALF_TO_US(val)) : "d"(a));
408 return val;
409,
410 __half result;
411 // Perform rounding to 11 bits of precision, convert value
412 // to float and call existing float to half conversion.
413 // By pre-rounding to 11 bits we avoid additional rounding
414 // in float to half conversion.
415 unsigned long long int absa;
416 unsigned long long int ua;
417 (void)memcpy(&ua, &a, sizeof(a));
418 absa = (ua & 0x7fffffffffffffffULL);
419 if ((absa >= 0x40f0000000000000ULL) || (absa <= 0x3e60000000000000ULL))
420 {
421 // |a| >= 2^16 or NaN or |a| <= 2^(-25)
422 // double-rounding is not a problem
423 result = __float2half(static_cast<float>(a));
424 }
425 else
426 {
427 // here 2^(-25) < |a| < 2^16
428 // prepare shifter value such that a + shifter
429 // done in double precision performs round-to-nearest-even
430 // and (a + shifter) - shifter results in a rounded to
431 // 11 bits of precision. Shifter needs to have exponent of
432 // a plus 53 - 11 = 42 and a leading bit in mantissa to guard
433 // against negative values.
434 // So need to have |a| capped to avoid overflow in exponent.
435 // For inputs that are smaller than half precision minnorm
436 // we prepare fixed shifter exponent.
437 unsigned long long shifterBits;
438 if (absa >= 0x3f10000000000000ULL)
439 { // Here if |a| >= 2^(-14)
440 // add 42 to exponent bits
441 shifterBits = (ua & 0x7ff0000000000000ULL) + 0x02A0000000000000ULL;
442 }
443 else
444 { // 2^(-25) < |a| < 2^(-14), potentially results in denormal
445 // set exponent bits to 42 - 14 + bias
446 shifterBits = 0x41B0000000000000ULL;
447 }
448 // set leading mantissa bit to protect against negative inputs
449 shifterBits |= 0x0008000000000000ULL;
450 double shifter;
451 (void)memcpy(&shifter, &shifterBits, sizeof(shifterBits));
452 double aShiftRound = a + shifter;
453
454 // Prevent the compiler from optimizing away a + shifter - shifter
455 // by doing intermediate memcopy and harmless bitwize operation
456 unsigned long long int aShiftRoundBits;
457 (void)memcpy(&aShiftRoundBits, &aShiftRound, sizeof(aShiftRound));
458
459 // the value is positive, so this operation doesn't change anything
460 aShiftRoundBits &= 0x7fffffffffffffffULL;
461
462 (void)memcpy(&aShiftRound, &aShiftRoundBits, sizeof(aShiftRound));
463
464 result = __float2half(static_cast<float>(aShiftRound - shifter));
465 }
466
467 return result;
468,
469 __half result;
470 /*
471 // Perform rounding to 11 bits of precision, convert value
472 // to float and call existing float to half conversion.
473 // By pre-rounding to 11 bits we avoid additional rounding
474 // in float to half conversion.
475 */
476 unsigned long long int absa;
477 unsigned long long int ua;
478 (void)std::memcpy(&ua, &a, sizeof(a));
479 absa = (ua & 0x7fffffffffffffffULL);
480 if ((absa >= 0x40f0000000000000ULL) || (absa <= 0x3e60000000000000ULL))
481 {
482 /*
483 // |a| >= 2^16 or NaN or |a| <= 2^(-25)
484 // double-rounding is not a problem
485 */
486 result = __float2half(static_cast<float>(a));
487 }
488 else
489 {
490 /*
491 // here 2^(-25) < |a| < 2^16
492 // prepare shifter value such that a + shifter
493 // done in double precision performs round-to-nearest-even
494 // and (a + shifter) - shifter results in a rounded to
495 // 11 bits of precision. Shifter needs to have exponent of
496 // a plus 53 - 11 = 42 and a leading bit in mantissa to guard
497 // against negative values.
498 // So need to have |a| capped to avoid overflow in exponent.
499 // For inputs that are smaller than half precision minnorm
500 // we prepare fixed shifter exponent.
501 */
502 unsigned long long shifterBits;
503 if (absa >= 0x3f10000000000000ULL)
504 {
505 /*
506 // Here if |a| >= 2^(-14)
507 // add 42 to exponent bits
508 */
509 shifterBits = (ua & 0x7ff0000000000000ULL) + 0x02A0000000000000ULL;
510 }
511 else
512 {
513 /*
514 // 2^(-25) < |a| < 2^(-14), potentially results in denormal
515 // set exponent bits to 42 - 14 + bias
516 */
517 shifterBits = 0x41B0000000000000ULL;
518 }
519 // set leading mantissa bit to protect against negative inputs
520 shifterBits |= 0x0008000000000000ULL;
521 double shifter;
522 (void)std::memcpy(&shifter, &shifterBits, sizeof(shifterBits));
523 double aShiftRound = a + shifter;
524
525 /*
526 // Prevent the compiler from optimizing away a + shifter - shifter
527 // by doing intermediate memcopy and harmless bitwize operation
528 */
529 unsigned long long int aShiftRoundBits;
530 (void)std::memcpy(&aShiftRoundBits, &aShiftRound, sizeof(aShiftRound));
531
532 // the value is positive, so this operation doesn't change anything
533 aShiftRoundBits &= 0x7fffffffffffffffULL;
534
535 (void)std::memcpy(&aShiftRound, &aShiftRoundBits, sizeof(aShiftRound));
536
537 result = __float2half(static_cast<float>(aShiftRound - shifter));
538 }
539
540 return result;
541)
542}
543__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half(const float a)
544{
545 __half val;
546NV_IF_ELSE_TARGET(NV_IS_DEVICE,
547 asm("{ cvt.rn.f16.f32 %0, %1;}\n" : "=h"(__HALF_TO_US(val)) : "f"(a));
548,
549 __half_raw r;
550 unsigned int sign = 0U;
551 unsigned int remainder = 0U;
552 r.x = __internal_float2half(a, sign, remainder);
553 if ((remainder > 0x80000000U) || ((remainder == 0x80000000U) && ((r.x & 0x1U) != 0U))) {
554 r.x++;
555 }
556 val = r;
557)
558 return val;
559}
560__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half_rn(const float a)
561{
562 __half val;
563NV_IF_ELSE_TARGET(NV_IS_DEVICE,
564 asm("{ cvt.rn.f16.f32 %0, %1;}\n" : "=h"(__HALF_TO_US(val)) : "f"(a));
565,
566 __half_raw r;
567 unsigned int sign = 0U;
568 unsigned int remainder = 0U;
569 r.x = __internal_float2half(a, sign, remainder);
570 if ((remainder > 0x80000000U) || ((remainder == 0x80000000U) && ((r.x & 0x1U) != 0U))) {
571 r.x++;
572 }
573 val = r;
574)
575 return val;
576}
577__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half_rz(const float a)
578{
579 __half val;
580NV_IF_ELSE_TARGET(NV_IS_DEVICE,
581 asm("{ cvt.rz.f16.f32 %0, %1;}\n" : "=h"(__HALF_TO_US(val)) : "f"(a));
582,
583 __half_raw r;
584 unsigned int sign = 0U;
585 unsigned int remainder = 0U;
586 r.x = __internal_float2half(a, sign, remainder);
587 val = r;
588)
589 return val;
590}
591__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half_rd(const float a)
592{
593 __half val;
594NV_IF_ELSE_TARGET(NV_IS_DEVICE,
595 asm("{ cvt.rm.f16.f32 %0, %1;}\n" : "=h"(__HALF_TO_US(val)) : "f"(a));
596,
597 __half_raw r;
598 unsigned int sign = 0U;
599 unsigned int remainder = 0U;
600 r.x = __internal_float2half(a, sign, remainder);
601 if ((remainder != 0U) && (sign != 0U)) {
602 r.x++;
603 }
604 val = r;
605)
606 return val;
607}
608__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half_ru(const float a)
609{
610 __half val;
611NV_IF_ELSE_TARGET(NV_IS_DEVICE,
612 asm("{ cvt.rp.f16.f32 %0, %1;}\n" : "=h"(__HALF_TO_US(val)) : "f"(a));
613,
614 __half_raw r;
615 unsigned int sign = 0U;
616 unsigned int remainder = 0U;
617 r.x = __internal_float2half(a, sign, remainder);
618 if ((remainder != 0U) && (sign == 0U)) {
619 r.x++;
620 }
621 val = r;
622)
623 return val;
624}
625__CUDA_HOSTDEVICE_FP16_DECL__ __half2 __float2half2_rn(const float a)
626{
627 __half2 val;
628NV_IF_ELSE_TARGET(NV_IS_DEVICE,
629 asm("{.reg .f16 low;\n"
630 " cvt.rn.f16.f32 low, %1;\n"
631 " mov.b32 %0, {low,low};}\n" : "=r"(__HALF2_TO_UI(val)) : "f"(a));
632,
633 val = __half2(__float2half_rn(a), __float2half_rn(a));
634)
635 return val;
636}
637
638#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
639__CUDA_FP16_DECL__ __half2 __internal_device_float2_to_half2_rn(const float a, const float b) {
640 __half2 val;
641NV_IF_ELSE_TARGET(NV_PROVIDES_SM_80,
642 asm("{ cvt.rn.f16x2.f32 %0, %2, %1; }\n"
643 : "=r"(__HALF2_TO_UI(val)) : "f"(a), "f"(b));
644,
645 asm("{.reg .f16 low,high;\n"
646 " cvt.rn.f16.f32 low, %1;\n"
647 " cvt.rn.f16.f32 high, %2;\n"
648 " mov.b32 %0, {low,high};}\n" : "=r"(__HALF2_TO_UI(val)) : "f"(a), "f"(b));
649)
650 return val;
651}
652
653#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
654
655__CUDA_HOSTDEVICE_FP16_DECL__ __half2 __floats2half2_rn(const float a, const float b)
656{
657 __half2 val;
658NV_IF_ELSE_TARGET(NV_IS_DEVICE,
659 val = __internal_device_float2_to_half2_rn(a,b);
660,
661 val = __half2(__float2half_rn(a), __float2half_rn(b));
662)
663 return val;
664}
665
666#ifndef __CUDACC_RTC__ /* no host functions in NVRTC mode */
667static inline float __internal_half2float(const unsigned short h)
668{
669 unsigned int sign = ((static_cast<unsigned int>(h) >> 15U) & 1U);
670 unsigned int exponent = ((static_cast<unsigned int>(h) >> 10U) & 0x1fU);
671 unsigned int mantissa = ((static_cast<unsigned int>(h) & 0x3ffU) << 13U);
672 float f;
673 if (exponent == 0x1fU) { /* NaN or Inf */
674 /* discard sign of a NaN */
675 sign = ((mantissa != 0U) ? (sign >> 1U) : sign);
676 mantissa = ((mantissa != 0U) ? 0x7fffffU : 0U);
677 exponent = 0xffU;
678 } else if (exponent == 0U) { /* Denorm or Zero */
679 if (mantissa != 0U) {
680 unsigned int msb;
681 exponent = 0x71U;
682 do {
683 msb = (mantissa & 0x400000U);
684 mantissa <<= 1U; /* normalize */
685 --exponent;
686 } while (msb == 0U);
687 mantissa &= 0x7fffffU; /* 1.mantissa is implicit */
688 }
689 } else {
690 exponent += 0x70U;
691 }
692 const unsigned int u = ((sign << 31U) | (exponent << 23U) | mantissa);
693#if defined(__CUDACC__)
694 (void)memcpy(&f, &u, sizeof(u));
695#else
696 (void)std::memcpy(&f, &u, sizeof(u));
697#endif
698 return f;
699}
700#endif /* !defined(__CUDACC_RTC__) */
701
702__CUDA_HOSTDEVICE_FP16_DECL__ float __half2float(const __half a)
703{
704 float val;
705NV_IF_ELSE_TARGET(NV_IS_DEVICE,
706 asm("{ cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(__HALF_TO_CUS(a)));
707,
708 val = __internal_half2float(static_cast<__half_raw>(a).x);
709)
710 return val;
711}
712__CUDA_HOSTDEVICE_FP16_DECL__ float __low2float(const __half2 a)
713{
714 float val;
715NV_IF_ELSE_TARGET(NV_IS_DEVICE,
716 asm("{.reg .f16 low,high;\n"
717 " mov.b32 {low,high},%1;\n"
718 " cvt.f32.f16 %0, low;}\n" : "=f"(val) : "r"(__HALF2_TO_CUI(a)));
719,
720 val = __internal_half2float(static_cast<__half2_raw>(a).x);
721)
722 return val;
723}
724__CUDA_HOSTDEVICE_FP16_DECL__ float __high2float(const __half2 a)
725{
726 float val;
727NV_IF_ELSE_TARGET(NV_IS_DEVICE,
728 asm("{.reg .f16 low,high;\n"
729 " mov.b32 {low,high},%1;\n"
730 " cvt.f32.f16 %0, high;}\n" : "=f"(val) : "r"(__HALF2_TO_CUI(a)));
731,
732 val = __internal_half2float(static_cast<__half2_raw>(a).y);
733)
734 return val;
735}
736
737__CUDA_HOSTDEVICE_FP16_DECL__ signed char __half2char_rz(const __half h)
738{
739 signed char i;
740NV_IF_ELSE_TARGET(NV_PROVIDES_SM_90,
741 unsigned int tmp;
742 asm("cvt.rzi.s8.f16 %0, %1;" : "=r"(tmp) : "h"(__HALF_TO_CUS(h)));
743 const unsigned char u = static_cast<unsigned char>(tmp);
744 i = static_cast<signed char>(u);
745,
746 const float f = __half2float(h);
747 const signed char max_val = (signed char)0x7fU;
748 const signed char min_val = (signed char)0x80U;
749 const unsigned short bits = static_cast<unsigned short>(static_cast<__half_raw>(h).x << 1U);
750 // saturation fixup
751 if (bits > (unsigned short)0xF800U) {
752 // NaN
753 i = 0;
754 } else if (f > static_cast<float>(max_val)) {
755 // saturate maximum
756 i = max_val;
757 } else if (f < static_cast<float>(min_val)) {
758 // saturate minimum
759 i = min_val;
760 } else {
761 // normal value, conversion is well-defined
762 i = static_cast<signed char>(f);
763 }
764)
765 return i;
766}
767
768__CUDA_HOSTDEVICE_FP16_DECL__ unsigned char __half2uchar_rz(const __half h)
769{
770 unsigned char i;
771NV_IF_ELSE_TARGET(NV_PROVIDES_SM_90,
772 unsigned int tmp;
773 asm("cvt.rzi.u8.f16 %0, %1;" : "=r"(tmp) : "h"(__HALF_TO_CUS(h)));
774 i = static_cast<unsigned char>(tmp);
775,
776 const float f = __half2float(h);
777 const unsigned char max_val = 0xffU;
778 const unsigned char min_val = 0U;
779 const unsigned short bits = static_cast<unsigned short>(static_cast<__half_raw>(h).x << 1U);
780 // saturation fixup
781 if (bits > (unsigned short)0xF800U) {
782 // NaN
783 i = 0U;
784 } else if (f > static_cast<float>(max_val)) {
785 // saturate maximum
786 i = max_val;
787 } else if (f < static_cast<float>(min_val)) {
788 // saturate minimum
789 i = min_val;
790 } else {
791 // normal value, conversion is well-defined
792 i = static_cast<unsigned char>(f);
793 }
794)
795 return i;
796}
797
798__CUDA_HOSTDEVICE_FP16_DECL__ short int __half2short_rz(const __half h)
799{
800 short int i;
801NV_IF_ELSE_TARGET(NV_IS_DEVICE,
802 asm("cvt.rzi.s16.f16 %0, %1;" : "=h"(i) : "h"(__HALF_TO_CUS(h)));
803,
804 const float f = __half2float(h);
805 const short int max_val = (short int)0x7fffU;
806 const short int min_val = (short int)0x8000U;
807 const unsigned short bits = static_cast<unsigned short>(static_cast<__half_raw>(h).x << 1U);
808 // saturation fixup
809 if (bits > (unsigned short)0xF800U) {
810 // NaN
811 i = 0;
812 } else if (f > static_cast<float>(max_val)) {
813 // saturate maximum
814 i = max_val;
815 } else if (f < static_cast<float>(min_val)) {
816 // saturate minimum
817 i = min_val;
818 } else {
819 // normal value, conversion is well-defined
820 i = static_cast<short int>(f);
821 }
822)
823 return i;
824}
825__CUDA_HOSTDEVICE_FP16_DECL__ unsigned short int __half2ushort_rz(const __half h)
826{
827 unsigned short int i;
828NV_IF_ELSE_TARGET(NV_IS_DEVICE,
829 asm("cvt.rzi.u16.f16 %0, %1;" : "=h"(i) : "h"(__HALF_TO_CUS(h)));
830,
831 const float f = __half2float(h);
832 const unsigned short int max_val = 0xffffU;
833 const unsigned short int min_val = 0U;
834 const unsigned short bits = static_cast<unsigned short>(static_cast<__half_raw>(h).x << 1U);
835 // saturation fixup
836 if (bits > (unsigned short)0xF800U) {
837 // NaN
838 i = 0U;
839 } else if (f > static_cast<float>(max_val)) {
840 // saturate maximum
841 i = max_val;
842 } else if (f < static_cast<float>(min_val)) {
843 // saturate minimum
844 i = min_val;
845 } else {
846 // normal value, conversion is well-defined
847 i = static_cast<unsigned short int>(f);
848 }
849)
850 return i;
851}
852__CUDA_HOSTDEVICE_FP16_DECL__ int __half2int_rz(const __half h)
853{
854 int i;
855NV_IF_ELSE_TARGET(NV_IS_DEVICE,
856 asm("cvt.rzi.s32.f16 %0, %1;" : "=r"(i) : "h"(__HALF_TO_CUS(h)));
857,
858 const float f = __half2float(h);
859 const int max_val = (int)0x7fffffffU;
860 const int min_val = (int)0x80000000U;
861 const unsigned short bits = static_cast<unsigned short>(static_cast<__half_raw>(h).x << 1U);
862 // saturation fixup
863 if (bits > (unsigned short)0xF800U) {
864 // NaN
865 i = 0;
866 } else if (f > static_cast<float>(max_val)) {
867 // saturate maximum
868 i = max_val;
869 } else if (f < static_cast<float>(min_val)) {
870 // saturate minimum
871 i = min_val;
872 } else {
873 // normal value, conversion is well-defined
874 i = static_cast<int>(f);
875 }
876)
877 return i;
878}
879__CUDA_HOSTDEVICE_FP16_DECL__ unsigned int __half2uint_rz(const __half h)
880{
881 unsigned int i;
882NV_IF_ELSE_TARGET(NV_IS_DEVICE,
883 asm("cvt.rzi.u32.f16 %0, %1;" : "=r"(i) : "h"(__HALF_TO_CUS(h)));
884,
885 const float f = __half2float(h);
886 const unsigned int max_val = 0xffffffffU;
887 const unsigned int min_val = 0U;
888 const unsigned short bits = static_cast<unsigned short>(static_cast<__half_raw>(h).x << 1U);
889 // saturation fixup
890 if (bits > (unsigned short)0xF800U) {
891 // NaN
892 i = 0U;
893 } else if (f > static_cast<float>(max_val)) {
894 // saturate maximum
895 i = max_val;
896 } else if (f < static_cast<float>(min_val)) {
897 // saturate minimum
898 i = min_val;
899 } else {
900 // normal value, conversion is well-defined
901 i = static_cast<unsigned int>(f);
902 }
903)
904 return i;
905}
906__CUDA_HOSTDEVICE_FP16_DECL__ long long int __half2ll_rz(const __half h)
907{
908 long long int i;
909NV_IF_ELSE_TARGET(NV_IS_DEVICE,
910 asm("cvt.rzi.s64.f16 %0, %1;" : "=l"(i) : "h"(__HALF_TO_CUS(h)));
911,
912 const float f = __half2float(h);
913 const long long int max_val = (long long int)0x7fffffffffffffffULL;
914 const long long int min_val = (long long int)0x8000000000000000ULL;
915 const unsigned short bits = static_cast<unsigned short>(static_cast<__half_raw>(h).x << 1U);
916 // saturation fixup
917 if (bits > (unsigned short)0xF800U) {
918 // NaN
919 i = min_val;
920 } else if (f > static_cast<float>(max_val)) {
921 // saturate maximum
922 i = max_val;
923 } else if (f < static_cast<float>(min_val)) {
924 // saturate minimum
925 i = min_val;
926 } else {
927 // normal value, conversion is well-defined
928 i = static_cast<long long int>(f);
929 }
930)
931 return i;
932}
933__CUDA_HOSTDEVICE_FP16_DECL__ unsigned long long int __half2ull_rz(const __half h)
934{
935 unsigned long long int i;
936NV_IF_ELSE_TARGET(NV_IS_DEVICE,
937 asm("cvt.rzi.u64.f16 %0, %1;" : "=l"(i) : "h"(__HALF_TO_CUS(h)));
938,
939 const float f = __half2float(h);
940 const unsigned long long int max_val = 0xffffffffffffffffULL;
941 const unsigned long long int min_val = 0ULL;
942 const unsigned short bits = static_cast<unsigned short>(static_cast<__half_raw>(h).x << 1U);
943 // saturation fixup
944 if (bits > (unsigned short)0xF800U) {
945 // NaN
946 i = 0x8000000000000000ULL;
947 } else if (f > static_cast<float>(max_val)) {
948 // saturate maximum
949 i = max_val;
950 } else if (f < static_cast<float>(min_val)) {
951 // saturate minimum
952 i = min_val;
953 } else {
954 // normal value, conversion is well-defined
955 i = static_cast<unsigned long long int>(f);
956 }
957)
958 return i;
959}
960/* CUDA vector-types compatible vector creation function (note returns __half2, not half2) */
961__CUDA_HOSTDEVICE_FP16_DECL__ __half2 make_half2(const __half x, const __half y)
962{
963 __half2 t; t.x = x; t.y = y; return t;
964}
965
966
967/* Definitions of intrinsics */
968__CUDA_HOSTDEVICE_FP16_DECL__ __half2 __float22half2_rn(const float2 a)
969{
970 const __half2 val = __floats2half2_rn(a.x, a.y);
971 return val;
972}
973__CUDA_HOSTDEVICE_FP16_DECL__ float2 __half22float2(const __half2 a)
974{
975 float hi_float;
976 float lo_float;
977NV_IF_ELSE_TARGET(NV_IS_DEVICE,
978 asm("{.reg .f16 low,high;\n"
979 " mov.b32 {low,high},%1;\n"
980 " cvt.f32.f16 %0, low;}\n" : "=f"(lo_float) : "r"(__HALF2_TO_CUI(a)));
981
982 asm("{.reg .f16 low,high;\n"
983 " mov.b32 {low,high},%1;\n"
984 " cvt.f32.f16 %0, high;}\n" : "=f"(hi_float) : "r"(__HALF2_TO_CUI(a)));
985,
986 lo_float = __internal_half2float(((__half2_raw)a).x);
987 hi_float = __internal_half2float(((__half2_raw)a).y);
988)
989 return make_float2(lo_float, hi_float);
990}
991#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
992__CUDA_FP16_DECL__ int __half2int_rn(const __half h)
993{
994 int i;
995 asm("cvt.rni.s32.f16 %0, %1;" : "=r"(i) : "h"(__HALF_TO_CUS(h)));
996 return i;
997}
998__CUDA_FP16_DECL__ int __half2int_rd(const __half h)
999{
1000 int i;
1001 asm("cvt.rmi.s32.f16 %0, %1;" : "=r"(i) : "h"(__HALF_TO_CUS(h)));
1002 return i;
1003}
1004__CUDA_FP16_DECL__ int __half2int_ru(const __half h)
1005{
1006 int i;
1007 asm("cvt.rpi.s32.f16 %0, %1;" : "=r"(i) : "h"(__HALF_TO_CUS(h)));
1008 return i;
1009}
1010#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
1011__CUDA_HOSTDEVICE_FP16_DECL__ __half __int2half_rn(const int i)
1012{
1013 __half h;
1014NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1015 asm("cvt.rn.f16.s32 %0, %1;" : "=h"(__HALF_TO_US(h)) : "r"(i));
1016,
1017 // double-rounding is not a problem here: if integer
1018 // has more than 24 bits, it is already too large to
1019 // be represented in half precision, and result will
1020 // be infinity.
1021 const float f = static_cast<float>(i);
1022 h = __float2half_rn(f);
1023)
1024 return h;
1025}
1026__CUDA_HOSTDEVICE_FP16_DECL__ __half __int2half_rz(const int i)
1027{
1028 __half h;
1029NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1030 asm("cvt.rz.f16.s32 %0, %1;" : "=h"(__HALF_TO_US(h)) : "r"(i));
1031,
1032 const float f = static_cast<float>(i);
1033 h = __float2half_rz(f);
1034)
1035 return h;
1036}
1037__CUDA_HOSTDEVICE_FP16_DECL__ __half __int2half_rd(const int i)
1038{
1039 __half h;
1040NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1041 asm("cvt.rm.f16.s32 %0, %1;" : "=h"(__HALF_TO_US(h)) : "r"(i));
1042,
1043 const float f = static_cast<float>(i);
1044 h = __float2half_rd(f);
1045)
1046 return h;
1047}
1048__CUDA_HOSTDEVICE_FP16_DECL__ __half __int2half_ru(const int i)
1049{
1050 __half h;
1051NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1052 asm("cvt.rp.f16.s32 %0, %1;" : "=h"(__HALF_TO_US(h)) : "r"(i));
1053,
1054 const float f = static_cast<float>(i);
1055 h = __float2half_ru(f);
1056)
1057 return h;
1058}
1059#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
1060__CUDA_FP16_DECL__ short int __half2short_rn(const __half h)
1061{
1062 short int i;
1063 asm("cvt.rni.s16.f16 %0, %1;" : "=h"(i) : "h"(__HALF_TO_CUS(h)));
1064 return i;
1065}
1066__CUDA_FP16_DECL__ short int __half2short_rd(const __half h)
1067{
1068 short int i;
1069 asm("cvt.rmi.s16.f16 %0, %1;" : "=h"(i) : "h"(__HALF_TO_CUS(h)));
1070 return i;
1071}
1072__CUDA_FP16_DECL__ short int __half2short_ru(const __half h)
1073{
1074 short int i;
1075 asm("cvt.rpi.s16.f16 %0, %1;" : "=h"(i) : "h"(__HALF_TO_CUS(h)));
1076 return i;
1077}
1078#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
1079__CUDA_HOSTDEVICE_FP16_DECL__ __half __short2half_rn(const short int i)
1080{
1081 __half h;
1082NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1083 asm("cvt.rn.f16.s16 %0, %1;" : "=h"(__HALF_TO_US(h)) : "h"(i));
1084,
1085 const float f = static_cast<float>(i);
1086 h = __float2half_rn(f);
1087)
1088 return h;
1089}
1090__CUDA_HOSTDEVICE_FP16_DECL__ __half __short2half_rz(const short int i)
1091{
1092 __half h;
1093NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1094 asm("cvt.rz.f16.s16 %0, %1;" : "=h"(__HALF_TO_US(h)) : "h"(i));
1095,
1096 const float f = static_cast<float>(i);
1097 h = __float2half_rz(f);
1098)
1099 return h;
1100}
1101__CUDA_HOSTDEVICE_FP16_DECL__ __half __short2half_rd(const short int i)
1102{
1103 __half h;
1104NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1105 asm("cvt.rm.f16.s16 %0, %1;" : "=h"(__HALF_TO_US(h)) : "h"(i));
1106,
1107 const float f = static_cast<float>(i);
1108 h = __float2half_rd(f);
1109)
1110 return h;
1111}
1112__CUDA_HOSTDEVICE_FP16_DECL__ __half __short2half_ru(const short int i)
1113{
1114 __half h;
1115NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1116 asm("cvt.rp.f16.s16 %0, %1;" : "=h"(__HALF_TO_US(h)) : "h"(i));
1117,
1118 const float f = static_cast<float>(i);
1119 h = __float2half_ru(f);
1120)
1121 return h;
1122}
1123#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
1124__CUDA_FP16_DECL__ unsigned int __half2uint_rn(const __half h)
1125{
1126 unsigned int i;
1127 asm("cvt.rni.u32.f16 %0, %1;" : "=r"(i) : "h"(__HALF_TO_CUS(h)));
1128 return i;
1129}
1130__CUDA_FP16_DECL__ unsigned int __half2uint_rd(const __half h)
1131{
1132 unsigned int i;
1133 asm("cvt.rmi.u32.f16 %0, %1;" : "=r"(i) : "h"(__HALF_TO_CUS(h)));
1134 return i;
1135}
1136__CUDA_FP16_DECL__ unsigned int __half2uint_ru(const __half h)
1137{
1138 unsigned int i;
1139 asm("cvt.rpi.u32.f16 %0, %1;" : "=r"(i) : "h"(__HALF_TO_CUS(h)));
1140 return i;
1141}
1142#endif /* defined(__CUDACC__) || defined(_NVHPC_CUDA) */
1143__CUDA_HOSTDEVICE_FP16_DECL__ __half __uint2half_rn(const unsigned int i)
1144{
1145 __half h;
1146NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1147 asm("cvt.rn.f16.u32 %0, %1;" : "=h"(__HALF_TO_US(h)) : "r"(i));
1148,
1149 // double-rounding is not a problem here: if integer
1150 // has more than 24 bits, it is already too large to
1151 // be represented in half precision, and result will
1152 // be infinity.
1153 const float f = static_cast<float>(i);
1154 h = __float2half_rn(f);
1155)
1156 return h;
1157}
1158__CUDA_HOSTDEVICE_FP16_DECL__ __half __uint2half_rz(const unsigned int i)
1159{
1160 __half h;
1161NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1162 asm("cvt.rz.f16.u32 %0, %1;" : "=h"(__HALF_TO_US(h)) : "r"(i));
1163,
1164 const float f = static_cast<float>(i);
1165 h = __float2half_rz(f);
1166)
1167 return h;
1168}
1169__CUDA_HOSTDEVICE_FP16_DECL__ __half __uint2half_rd(const unsigned int i)
1170{
1171 __half h;
1172NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1173 asm("cvt.rm.f16.u32 %0, %1;" : "=h"(__HALF_TO_US(h)) : "r"(i));
1174,
1175 const float f = static_cast<float>(i);
1176 h = __float2half_rd(f);
1177)
1178 return h;
1179}
1180__CUDA_HOSTDEVICE_FP16_DECL__ __half __uint2half_ru(const unsigned int i)
1181{
1182 __half h;
1183NV_IF_ELSE_TARGET(NV_IS_DEVICE,
1184 asm("cvt.rp.f16.u32 %0, %1;" : "=h"(__HALF_TO_US(h)) : "r"(i));
1185,
1186 const float f = static_cast<float>(i);
1187 h = __float2half_ru(f);
1188)
1189 return h;
1190}
1191#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
1192__CUDA_FP16_DECL__ unsigned short int __half2ushort_rn(const __half h)
1193{
1194 unsigned short int i;
1195 asm("cvt.rni.u16.f16 %0, %1;" : "=h"(i) : "h"(__HALF_TO_CUS(h)));
1196 return i;
1197}
1198__CUDA_FP16_DECL__ unsigned short int __half2ushort_rd(const __half h)
1199{
1200 unsigned short int i;