|
|
CUDA Math API Reference Manual
Release 12.9
NVIDIA Corporation
May 31, 2025
Contents
1
FP4 Intrinsics
3
1.1
C++ struct for handling fp4 data type of e2m1 kind
3
1.2
C++ struct for handling vector type of four fp4 values of e2m1 kind
3
1.3
C++ struct for handling vector type of two fp4 values of e2m1 kind
4
1.4
FP4 Conversion and Data Movement
4
1.4.1
Enumerations
7
1.4.2
Functions
7
1.4.3
Typedefs
10
2
FP6 Intrinsics
13
2.1
C++ struct for handling fp6 data type of e2m3 kind
13
2.2
C++ struct for handling fp6 data type of e3m2 kind
13
2.3
C++ struct for handling vector type of four fp6 values of e2m3 kind
14
2.4
C++ struct for handling vector type of four fp6 values of e3m2 kind
14
2.5
C++ struct for handling vector type of two fp6 values of e2m3 kind
14
2.6
C++ struct for handling vector type of two fp6 values of e3m2 kind
14
2.7
FP6 Conversion and Data Movement
15
2.7.1
Enumerations
19
2.7.2
Functions
19
2.7.3
Typedefs
22
3
FP8 Intrinsics
25
3.1
C++ struct for handling fp8 data type of e4m3 kind
25
3.2
C++ struct for handling fp8 data type of e5m2 kind
25
3.3
C++ struct for handling vector type of four fp8 values of e4m3 kind
26
3.4
C++ struct for handling vector type of four fp8 values of e5m2 kind
26
3.5
C++ struct for handling vector type of four scale factors of e8m0 kind.
26
3.6
C++ struct for handling vector type of two fp8 values of e4m3 kind
26
3.7
C++ struct for handling vector type of two fp8 values of e5m2 kind
27
3.8
C++ struct for handling vector type of two scale factors of e8m0 kind.
27
3.9
FP8 Conversion and Data Movement
27
3.9.1
Enumerations
36
3.9.2
Functions
37
3.9.3
Typedefs
42
4
Half Precision Intrinsics
43
4.1
Half Arithmetic Constants
43
4.1.1
Macros
44
4.2
Half Arithmetic Functions
44
4.2.1
Functions
46
4.3
Half Comparison Functions
53
4.3.1
Functions
54
4.4
Half Math Functions
60
i
4.4.1
Functions
61
4.5
Half Precision Conversion and Data Movement
66
4.5.1
Functions
74
4.6
Half2 Arithmetic Functions
104
4.6.1
Functions
105
4.7
Half2 Comparison Functions
112
4.7.1
Functions
115
4.8
Half2 Math Functions
129
4.8.1
Functions
130
4.9
Typedefs
136
5
Bfloat16 Precision Intrinsics
137
5.1
Bfloat16 Arithmetic Constants
137
5.1.1
Macros
138
5.2
Bfloat16 Arithmetic Functions
139
5.2.1
Functions
140
5.3
Bfloat16 Comparison Functions
146
5.3.1
Functions
147
5.4
Bfloat16 Math Functions
153
5.4.1
Functions
154
5.5
Bfloat16 Precision Conversion and Data Movement
158
5.5.1
Functions
166
5.6
Bfloat162 Arithmetic Functions
194
5.6.1
Functions
196
5.7
Bfloat162 Comparison Functions
202
5.7.1
Functions
205
5.8
Bfloat162 Math Functions
218
5.8.1
Functions
219
5.9
Typedefs
224
6
Single Precision Mathematical Functions
225
6.1
Functions
229
7
Single Precision Intrinsics
261
7.1
Functions
264
8
Double Precision Mathematical Functions
285
8.1
Functions
290
9
Double Precision Intrinsics
321
9.1
Functions
322
10 FP128 Quad Precision Mathematical Functions
337
10.1
Functions
340
11 Type Casting Intrinsics
359
11.1
Functions
362
12 Integer Mathematical Functions
377
12.1
Functions
379
13 Integer Intrinsics
383
13.1
Functions
385
14 SIMD Intrinsics
393
14.1
Functions
399
ii
15 Structs
419
15.1
__half
419
15.2
__half2
423
15.3
__half2_raw
424
15.4
__half_raw
425
15.5
__nv_bfloat16
425
15.6
__nv_bfloat162
428
15.7
__nv_bfloat162_raw
429
15.8
__nv_bfloat16_raw . .
430
15.9
__nv_fp4_e2m1
430
15.10
__nv_fp4x2_e2m1
432
15.11
__nv_fp4x4_e2m1
433
15.12
__nv_fp6_e2m3
433
15.13
__nv_fp6_e3m2
435
15.14
__nv_fp6x2_e2m3
436
15.15
__nv_fp6x2_e3m2
437
15.16
__nv_fp6x4_e2m3
438
15.17
__nv_fp6x4_e3m2
438
15.18
__nv_fp8_e4m3
439
15.19
__nv_fp8_e5m2
442
15.20
__nv_fp8_e8m0
445
15.21
__nv_fp8x2_e4m3
448
15.22
__nv_fp8x2_e5m2
449
15.23
__nv_fp8x2_e8m0
450
15.24
__nv_fp8x4_e4m3
451
15.25
__nv_fp8x4_e5m2
451
15.26
__nv_fp8x4_e8m0
452
16 Notices
455
16.1
Notice
455
16.2
OpenCL
456
16.3
Trademarks
456
iii
iv
CUDA mathematical functions are always available in device code.
Host implementations of the common mathematical functions are mapped in a platform-specific
way to standard math library functions, provided by the host compiler and respective host libm
where available.
Some functions, not available with the host compilers, are implemented in
crt/math_functions.hpp header file. For example, see erfinv(). Other, less common functions, like rhy-
pot(), cyl_bessel_i0() are only available in device code.
CUDA Math device functions are no-throw for well-formed CUDA programs.
Note that many floating-point and integer functions names are overloaded for different argument
types. For example, the log() function has the following prototypes:
double log(double x);
float log(float x);
float logf(float x);
Note also that due to implementation constraints, certain math functions from std:: namespace may
be callable in device code even via explicitly qualified std:: names. However, such use is discouraged,
since this capability is unsupported, unverified, undocumented, not portable, and may change without
notice.
Contents
1
2
Contents
Chapter 1. FP4 Intrinsics
This section describes fp4 intrinsic functions.
To use these functions, include the header file cuda_fp4.h in your program.
The following macros are available to help users selectively enable/disable various definitions present
in the header file:
▶ __CUDA_NO_FP4_CONVERSIONS__ - If defined, this macro will prevent any use of the C++ type
conversions (converting constructors and conversion operators) defined in the header.
▶ __CUDA_NO_FP4_CONVERSION_OPERATORS__ - If defined, this macro will prevent any use of the
C++ conversion operators from fp4 to other types.
Note: Most of the operations defined here benefit from native HW support when compiled for specific
GPU targets (e.g. devices of compute capability 10.0a), other targets use emulation path.
1.1. C++ struct for handling fp4 data type of
e2m1 kind.
Structs
__nv_fp4_e2m1
__nv_fp4_e2m1 datatype
1.2. C++ struct for handling vector type of four
fp4 values of e2m1 kind.
Structs
__nv_fp4x4_e2m1
__nv_fp4x4_e2m1 datatype
3
1.3. C++ struct for handling vector type of two
fp4 values of e2m1 kind.
Structs
__nv_fp4x2_e2m1
__nv_fp4x2_e2m1 datatype
1.4. FP4 Conversion and Data Movement
To use these functions, include the header file cuda_fp4.h in your program.
Enumerations
__nv_fp4_interpretation_t
Enumerates the possible interpretations of the 4-bit values when referring to them as fp4 types.
Functions
__host__
__device__
__nv_fp4x2_storage_t
__nv_cvt_bfloat16raw2_to_fp4x2(const
__nv_bfloat162_raw x, const __nv_fp4_interpretation_t fp4_interpretation, const enum
cudaRoundMode rounding)
Converts input vector of two nv_bfloat16 precision numbers packed in __nv_bfloat162_raw
x into a vector of two values of fp4 type of the requested kind using specified rounding mode
and saturating the out-of-range values.
__host__ __device__ __nv_fp4_storage_t __nv_cvt_bfloat16raw_to_fp4(const __nv_bfloat16_raw x,
const __nv_fp4_interpretation_t fp4_interpretation, const enum cudaRoundMode rounding)
Converts input nv_bfloat16 precision x to fp4 type of the requested kind using specified
rounding mode and saturating the out-of-range values.
__host__ __device__ __nv_fp4x2_storage_t __nv_cvt_double2_to_fp4x2(const double2 x, const
__nv_fp4_interpretation_t fp4_interpretation, const enum cudaRoundMode rounding)
Converts input vector of two double precision numbers packed in double2 x into a vector of
two values of fp4 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
__host__
__device__
__nv_fp4_storage_t
__nv_cvt_double_to_fp4(const double x,
const
__nv_fp4_interpretation_t fp4_interpretation, const enum cudaRoundMode rounding)
Converts input double precision x to fp4 type of the requested kind using specified rounding
mode and saturating the out-of-range values.
__host__
__device__
__nv_fp4x2_storage_t
__nv_cvt_float2_to_fp4x2(const float2 x, const
__nv_fp4_interpretation_t fp4_interpretation, const enum cudaRoundMode rounding)
Converts input vector of two single precision numbers packed in float2 x into a vector of
two values of fp4 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
4
Chapter 1. FP4 Intrinsics
__host__
__device__
__nv_fp4_storage_t
__nv_cvt_float_to_fp4(const
float
x,
const
__nv_fp4_interpretation_t fp4_interpretation, const enum cudaRoundMode rounding)
Converts input single precision x to fp4 type of the requested kind using specified rounding
mode and saturating the out-of-range values.
__host__
__device__
__half_raw
__nv_cvt_fp4_to_halfraw(const
__nv_fp4_storage_t x, const
__nv_fp4_interpretation_t fp4_interpretation)
Converts input fp4 x of the specified kind to half precision.
__host__ __device__ __half2_raw __nv_cvt_fp4x2_to_halfraw2(const __nv_fp4x2_storage_t x, const
__nv_fp4_interpretation_t fp4_interpretation)
Converts input vector of two fp4 values of the specified kind to a vector of two half precision
values packed in __half2_raw structure.
__host__ __device__ __nv_fp4x2_storage_t __nv_cvt_halfraw2_to_fp4x2(const __half2_raw x, const
__nv_fp4_interpretation_t fp4_interpretation, const enum cudaRoundMode rounding)
Converts input vector of two half precision numbers packed in __half2_raw x into a vector of
two values of fp4 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
__host__
__device__
__nv_fp4_storage_t
__nv_cvt_halfraw_to_fp4(const
__half_raw x, const
__nv_fp4_interpretation_t fp4_interpretation, const enum cudaRoundMode rounding)
Converts input half precision x to fp4 type of the requested kind using specified rounding
mode and saturating the out-of-range values.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1()
Constructor by default.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const unsigned long int val)
Constructor from unsigned long int data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for out-of-range values
and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const long int val)
Constructor from long int data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior for out-of-range values
and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const int val)
Constructor from int data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const unsigned short int val)
Constructor from unsigned short int data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const long long int val)
Constructor from long long int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const short int val)
Constructor from short int data type.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for out-of-
range values and cudaRoundNearest rounding mode.
1.4. FP4 Conversion and Data Movement
5
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const unsigned int val)
Constructor from unsigned int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ __nv_fp4_e2m1::__nv_fp4_e2m1(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for out-of-range values
and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp4x2_e2m1::__nv_fp4x2_e2m1(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp4x2_e2m1::__nv_fp4x2_e2m1(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ __nv_fp4x2_e2m1::__nv_fp4x2_e2m1(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp4x2_e2m1::__nv_fp4x2_e2m1(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp4x2_e2m1::__nv_fp4x2_e2m1()
Constructor by default.
__host__ __device__ __nv_fp4x4_e2m1::__nv_fp4x4_e2m1()
Constructor by default.
__host__
__device__
__nv_fp4x4_e2m1::__nv_fp4x4_e2m1(const
__nv_bfloat162 flo,
const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
__host__ __device__ __nv_fp4x4_e2m1::__nv_fp4x4_e2m1(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ __nv_fp4x4_e2m1::__nv_fp4x4_e2m1(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp4x4_e2m1::__nv_fp4x4_e2m1(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior for
out-of-range values.
6
Chapter 1. FP4 Intrinsics
Typedefs
__nv_fp4_storage_t
8-bit unsigned integer type abstraction used for fp4 floating-point numbers storage.
__nv_fp4x2_storage_t
8-bit unsigned integer type abstraction used for storage of pairs of fp4 floating-point num-
bers.
__nv_fp4x4_storage_t
16-bit unsigned integer type abstraction used for storage of tetrads of fp4 floating-point
numbers.
1.4.1. Enumerations
enum __nv_fp4_interpretation_t
Enumerates the possible interpretations of the 4-bit values when referring to them as fp4 types.
Values:
enumerator __NV_E2M1
Stands for fp4 numbers of e2m1 kind.
1.4.2. Functions
__host__ __device__ __nv_fp4x2_storage_t __nv_cvt_bfloat16raw2_to_fp4x2(const
__nv_bfloat162_raw
x, const
__nv_fp4_interpretation_t
fp4_interpretation,
const enum
cudaRoundMode
rounding)
Converts input vector of two nv_bfloat16 precision numbers packed in __nv_bfloat162_raw
x into a vector of two values of fp4 type of the requested kind using specified rounding mode
and saturating the out-of-range values.
Converts input vector x to a vector of two fp4 values of the kind specified by
fp4_interpretation parameter, using rounding mode specified by rounding parame-
ter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result
in positive MAXNORM.
Returns
▶ The __nv_fp4x2_storage_t value holds the result of conversion.
1.4. FP4 Conversion and Data Movement
7
__host__ __device__ __nv_fp4_storage_t __nv_cvt_bfloat16raw_to_fp4(const
__nv_bfloat16_raw x,
const
__nv_fp4_interpretation_t
fp4_interpretation,
const enum
cudaRoundMode
rounding)
Converts input nv_bfloat16 precision x to fp4 type of the requested kind using specified
rounding mode and saturating the out-of-range values.
Converts input x to fp4 type of the kind specified by fp4_interpretation parameter, us-
ing rounding mode specified by rounding parameter. Large out-of-range values saturate to
MAXNORM of the same sign. NaN input values result in positive MAXNORM.
Returns
▶ The __nv_fp4_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp4x2_storage_t __nv_cvt_double2_to_fp4x2(const double2 x, const
__nv_fp4_interpretation_t
fp4_interpretation,
const enum
cudaRoundMode
rounding)
Converts input vector of two double precision numbers packed in double2 x into a vector of
two values of fp4 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
Converts input vector x to a vector of two fp4 values of the kind specified by
fp4_interpretation parameter, using rounding mode specified by rounding parame-
ter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result
in positive MAXNORM.
Returns
▶ The __nv_fp4x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp4_storage_t __nv_cvt_double_to_fp4(const double x, const
__nv_fp4_interpretation_t
fp4_interpretation, const
enum cudaRoundMode
rounding)
Converts input double precision x to fp4 type of the requested kind using specified rounding
mode and saturating the out-of-range values.
Converts input x to fp4 type of the kind specified by fp4_interpretation parameter, us-
ing rounding mode specified by rounding parameter. Large out-of-range values saturate to
MAXNORM of the same sign. NaN input values result in positive MAXNORM.
Returns
▶ The __nv_fp4_storage_t value holds the result of conversion.
8
Chapter 1. FP4 Intrinsics
__host__ __device__ __nv_fp4x2_storage_t __nv_cvt_float2_to_fp4x2(const float2 x, const
__nv_fp4_interpretation_t
fp4_interpretation, const
enum cudaRoundMode
rounding)
Converts input vector of two single precision numbers packed in float2 x into a vector of
two values of fp4 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
Converts input vector x to a vector of two fp4 values of the kind specified by
fp4_interpretation parameter, using rounding mode specified by rounding parame-
ter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result
in positive MAXNORM.
Returns
▶ The __nv_fp4x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp4_storage_t __nv_cvt_float_to_fp4(const float x, const
__nv_fp4_interpretation_t
fp4_interpretation, const enum
cudaRoundMode rounding)
Converts input single precision x to fp4 type of the requested kind using specified rounding
mode and saturating the out-of-range values.
Converts input x to fp4 type of the kind specified by fp4_interpretation parameter, us-
ing rounding mode specified by rounding parameter. Large out-of-range values saturate to
MAXNORM of the same sign. NaN input values result in positive MAXNORM.
Returns
▶ The __nv_fp4_storage_t value holds the result of conversion.
__host__ __device__ __half_raw __nv_cvt_fp4_to_halfraw(const __nv_fp4_storage_t x, const
__nv_fp4_interpretation_t
fp4_interpretation)
Converts input fp4 x of the specified kind to half precision.
Converts input x of fp4 type of the kind specified by fp4_interpretation parameter to half
precision.
Returns
▶ The __half_raw value holds the result of conversion.
__host__ __device__ __half2_raw __nv_cvt_fp4x2_to_halfraw2(const __nv_fp4x2_storage_t x,
const __nv_fp4_interpretation_t
fp4_interpretation)
Converts input vector of two fp4 values of the specified kind to a vector of two half precision
values packed in __half2_raw structure.
Converts input vector x of fp4 type of the kind specified by fp4_interpretation parameter
to a vector of two half precision values and returns as __half2_raw structure.
Returns
1.4. FP4 Conversion and Data Movement
9
▶ The __half2_raw value holds the result of conversion.
__host__ __device__ __nv_fp4x2_storage_t __nv_cvt_halfraw2_to_fp4x2(const __half2_raw x,
const
__nv_fp4_interpretation_t
fp4_interpretation,
const enum
cudaRoundMode
rounding)
Converts input vector of two half precision numbers packed in __half2_raw x into a vector of
two values of fp4 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
Converts input vector x to a vector of two fp4 values of the kind specified by
fp4_interpretation parameter, using rounding mode specified by rounding parame-
ter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result
in positive MAXNORM.
Returns
▶ The __nv_fp4x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp4_storage_t __nv_cvt_halfraw_to_fp4(const __half_raw x, const
__nv_fp4_interpretation_t
fp4_interpretation, const
enum cudaRoundMode
rounding)
Converts input half precision x to fp4 type of the requested kind using specified rounding mode
and saturating the out-of-range values.
Converts input x to fp4 type of the kind specified by fp4_interpretation parameter, us-
ing rounding mode specified by rounding parameter. Large out-of-range values saturate to
MAXNORM of the same sign. NaN input values result in positive MAXNORM.
Returns
▶ The __nv_fp4_storage_t value holds the result of conversion.
1.4.3. Typedefs
typedef __nv_fp8_storage_t __nv_fp4_storage_t
8-bit unsigned integer type abstraction used for fp4 floating-point numbers storage.
typedef __nv_fp8_storage_t __nv_fp4x2_storage_t
8-bit unsigned integer type abstraction used for storage of pairs of fp4 floating-point num-
bers.
typedef __nv_fp8x2_storage_t __nv_fp4x4_storage_t
16-bit unsigned integer type abstraction used for storage of tetrads of fp4 floating-point
numbers.
10
Chapter 1. FP4 Intrinsics
Groups
C++ struct for handling fp4 data type of e2m1 kind.
C++ struct for handling vector type of four fp4 values of e2m1 kind.
C++ struct for handling vector type of two fp4 values of e2m1 kind.
FP4 Conversion and Data Movement
To use these functions, include the header file cuda_fp4.h in your program.
1.4. FP4 Conversion and Data Movement
11
12
Chapter 1. FP4 Intrinsics
Chapter 2. FP6 Intrinsics
This section describes fp6 intrinsic functions.
To use these functions, include the header file cuda_fp6.h in your program.
The following macros are available to help users selectively enable/disable various definitions present
in the header file:
▶ __CUDA_NO_FP6_CONVERSIONS__ - If defined, this macro will prevent any use of the C++ type
conversions (converting constructors and conversion operators) defined in the header.
▶ __CUDA_NO_FP6_CONVERSION_OPERATORS__ - If defined, this macro will prevent any use of the
C++ conversion operators from fp6 to other types.
Note: Most of the operations defined here benefit from native HW support when compiled for specific
GPU targets (e.g. devices of compute capability 10.0a), other targets use emulation path.
2.1. C++ struct for handling fp6 data type of
e2m3 kind.
Structs
__nv_fp6_e2m3
__nv_fp6_e2m3 datatype
2.2. C++ struct for handling fp6 data type of
e3m2 kind.
Structs
__nv_fp6_e3m2
__nv_fp6_e3m2 datatype
13
2.3. C++ struct for handling vector type of four
fp6 values of e2m3 kind.
Structs
__nv_fp6x4_e2m3
__nv_fp6x4_e2m3 datatype
2.4. C++ struct for handling vector type of four
fp6 values of e3m2 kind.
Structs
__nv_fp6x4_e3m2
__nv_fp6x4_e3m2 datatype
2.5. C++ struct for handling vector type of two
fp6 values of e2m3 kind.
Structs
__nv_fp6x2_e2m3
__nv_fp6x2_e2m3 datatype
2.6. C++ struct for handling vector type of two
fp6 values of e3m2 kind.
Structs
__nv_fp6x2_e3m2
__nv_fp6x2_e3m2 datatype
14
Chapter 2. FP6 Intrinsics
2.7. FP6 Conversion and Data Movement
To use these functions, include the header file cuda_fp6.h in your program.
Enumerations
__nv_fp6_interpretation_t
Enumerates the possible interpretations of the 8-bit values when referring to them as fp6 types.
Functions
__host__
__device__
__nv_fp6x2_storage_t
__nv_cvt_bfloat16raw2_to_fp6x2(const
__nv_bfloat162_raw x, const __nv_fp6_interpretation_t fp6_interpretation, const enum
cudaRoundMode rounding)
Converts input vector of two nv_bfloat16 precision numbers packed in __nv_bfloat162_raw
x into a vector of two values of fp6 type of the requested kind using specified rounding mode
and saturating the out-of-range values.
__host__ __device__ __nv_fp6_storage_t __nv_cvt_bfloat16raw_to_fp6(const __nv_bfloat16_raw x,
const __nv_fp6_interpretation_t fp6_interpretation, const enum cudaRoundMode rounding)
Converts input nv_bfloat16 precision x to fp6 type of the requested kind using specified
rounding mode and saturating the out-of-range values.
__host__ __device__ __nv_fp6x2_storage_t __nv_cvt_double2_to_fp6x2(const double2 x, const
__nv_fp6_interpretation_t fp6_interpretation, const enum cudaRoundMode rounding)
Converts input vector of two double precision numbers packed in double2 x into a vector of
two values of fp6 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
__host__
__device__
__nv_fp6_storage_t
__nv_cvt_double_to_fp6(const double x,
const
__nv_fp6_interpretation_t fp6_interpretation, const enum cudaRoundMode rounding)
Converts input double precision x to fp6 type of the requested kind using specified rounding
mode and saturating the out-of-range values.
__host__
__device__
__nv_fp6x2_storage_t
__nv_cvt_float2_to_fp6x2(const float2 x, const
__nv_fp6_interpretation_t fp6_interpretation, const enum cudaRoundMode rounding)
Converts input vector of two single precision numbers packed in float2 x into a vector of
two values of fp6 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
__host__
__device__
__nv_fp6_storage_t
__nv_cvt_float_to_fp6(const
float
x,
const
__nv_fp6_interpretation_t fp6_interpretation, const enum cudaRoundMode rounding)
Converts input single precision x to fp6 type of the requested kind using specified rounding
mode and saturating the out-of-range values.
__host__
__device__
__half_raw
__nv_cvt_fp6_to_halfraw(const
__nv_fp6_storage_t x, const
__nv_fp6_interpretation_t fp6_interpretation)
Converts input fp6 x of the specified kind to half precision.
__host__ __device__ __half2_raw __nv_cvt_fp6x2_to_halfraw2(const __nv_fp6x2_storage_t x, const
__nv_fp6_interpretation_t fp6_interpretation)
Converts input vector of two fp6 values of the specified kind to a vector of two half precision
values packed in __half2_raw structure.
2.7. FP6 Conversion and Data Movement
15
__host__ __device__ __nv_fp6x2_storage_t __nv_cvt_halfraw2_to_fp6x2(const __half2_raw x, const
__nv_fp6_interpretation_t fp6_interpretation, const enum cudaRoundMode rounding)
Converts input vector of two half precision numbers packed in __half2_raw x into a vector of
two values of fp6 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
__host__
__device__
__nv_fp6_storage_t
__nv_cvt_halfraw_to_fp6(const
__half_raw x, const
__nv_fp6_interpretation_t fp6_interpretation, const enum cudaRoundMode rounding)
Converts input half precision x to fp6 type of the requested kind using specified rounding
mode and saturating the out-of-range values.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3()
Constructor by default.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior for out-of-range values
and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const long long int val)
Constructor from long long int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const long int val)
Constructor from long int data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const unsigned short int val)
Constructor from unsigned short int data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for out-of-range values
and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const short int val)
Constructor from short int data type.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for out-of-
range values and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const int val)
Constructor from int data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for out-of-range values
and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const unsigned int val)
Constructor from unsigned int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp6_e2m3::__nv_fp6_e2m3(const unsigned long int val)
Constructor from unsigned long int data type, relies on __NV_SATFINITE behavior for out-
of-range values.
16
Chapter 2. FP6 Intrinsics
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const long long int val)
Constructor from long long int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for out-of-range values
and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const int val)
Constructor from int data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior for out-of-range values
and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const short int val)
Constructor from short int data type.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const long int val)
Constructor from long int data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const unsigned short int val)
Constructor from unsigned short int data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for out-of-
range values and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for out-of-range values
and cudaRoundNearest rounding mode.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const unsigned int val)
Constructor from unsigned int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2()
Constructor by default.
__host__ __device__ __nv_fp6_e3m2::__nv_fp6_e3m2(const unsigned long int val)
Constructor from unsigned long int data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ __nv_fp6x2_e2m3::__nv_fp6x2_e2m3(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ __nv_fp6x2_e2m3::__nv_fp6x2_e2m3(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp6x2_e2m3::__nv_fp6x2_e2m3(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
2.7. FP6 Conversion and Data Movement
17
__host__ __device__ __nv_fp6x2_e2m3::__nv_fp6x2_e2m3(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp6x2_e2m3::__nv_fp6x2_e2m3()
Constructor by default.
__host__ __device__ __nv_fp6x2_e3m2::__nv_fp6x2_e3m2(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp6x2_e3m2::__nv_fp6x2_e3m2(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp6x2_e3m2::__nv_fp6x2_e3m2(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ __nv_fp6x2_e3m2::__nv_fp6x2_e3m2(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp6x2_e3m2::__nv_fp6x2_e3m2()
Constructor by default.
__host__ __device__ __nv_fp6x4_e2m3::__nv_fp6x4_e2m3(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ __nv_fp6x4_e2m3::__nv_fp6x4_e2m3(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp6x4_e2m3::__nv_fp6x4_e2m3(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__
__device__
__nv_fp6x4_e2m3::__nv_fp6x4_e2m3(const
__nv_bfloat162 flo,
const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
__host__ __device__ __nv_fp6x4_e2m3::__nv_fp6x4_e2m3()
Constructor by default.
__host__ __device__ __nv_fp6x4_e3m2::__nv_fp6x4_e3m2()
Constructor by default.
__host__ __device__ __nv_fp6x4_e3m2::__nv_fp6x4_e3m2(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ __nv_fp6x4_e3m2::__nv_fp6x4_e3m2(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__
__device__
__nv_fp6x4_e3m2::__nv_fp6x4_e3m2(const
__nv_bfloat162 flo,
const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
18
Chapter 2. FP6 Intrinsics
__host__ __device__ __nv_fp6x4_e3m2::__nv_fp6x4_e3m2(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
Typedefs
__nv_fp6_storage_t
8-bit unsigned integer type abstraction used for fp6 floating-point numbers storage.
__nv_fp6x2_storage_t
16-bit unsigned integer type abstraction used for storage of pairs of fp6 floating-point num-
bers.
__nv_fp6x4_storage_t
32-bit unsigned integer type abstraction used for storage of tetrads of fp6 floating-point
numbers.
2.7.1. Enumerations
enum __nv_fp6_interpretation_t
Enumerates the possible interpretations of the 8-bit values when referring to them as fp6 types.
Values:
enumerator __NV_E2M3
Stands for fp6 numbers of e2m3 kind.
enumerator __NV_E3M2
Stands for fp6 numbers of e3m2 kind.
2.7.2. Functions
__host__ __device__ __nv_fp6x2_storage_t __nv_cvt_bfloat16raw2_to_fp6x2(const
__nv_bfloat162_raw
x, const
__nv_fp6_interpretation_t
fp6_interpretation,
const enum
cudaRoundMode
rounding)
Converts input vector of two nv_bfloat16 precision numbers packed in __nv_bfloat162_raw
x into a vector of two values of fp6 type of the requested kind using specified rounding mode
and saturating the out-of-range values.
Converts input vector x to a vector of two fp6 values of the kind specified by
fp6_interpretation parameter, using rounding mode specified by rounding parame-
ter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result
in positive MAXNORM.
2.7. FP6 Conversion and Data Movement
19
Returns
▶ The __nv_fp6x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp6_storage_t __nv_cvt_bfloat16raw_to_fp6(const
__nv_bfloat16_raw x,
const
__nv_fp6_interpretation_t
fp6_interpretation,
const enum
cudaRoundMode
rounding)
Converts input nv_bfloat16 precision x to fp6 type of the requested kind using specified
rounding mode and saturating the out-of-range values.
Converts input x to fp6 type of the kind specified by fp6_interpretation parameter, us-
ing rounding mode specified by rounding parameter. Large out-of-range values saturate to
MAXNORM of the same sign. NaN input values result in positive MAXNORM.
Returns
▶ The __nv_fp6_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp6x2_storage_t __nv_cvt_double2_to_fp6x2(const double2 x, const
__nv_fp6_interpretation_t
fp6_interpretation,
const enum
cudaRoundMode
rounding)
Converts input vector of two double precision numbers packed in double2 x into a vector of
two values of fp6 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
Converts input vector x to a vector of two fp6 values of the kind specified by
fp6_interpretation parameter, using rounding mode specified by rounding parame-
ter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result
in positive MAXNORM.
Returns
▶ The __nv_fp6x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp6_storage_t __nv_cvt_double_to_fp6(const double x, const
__nv_fp6_interpretation_t
fp6_interpretation, const
enum cudaRoundMode
rounding)
Converts input double precision x to fp6 type of the requested kind using specified rounding
mode and saturating the out-of-range values.
Converts input x to fp6 type of the kind specified by fp6_interpretation parameter, us-
ing rounding mode specified by rounding parameter. Large out-of-range values saturate to
MAXNORM of the same sign. NaN input values result in positive MAXNORM.
Returns
20
Chapter 2. FP6 Intrinsics
▶ The __nv_fp6_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp6x2_storage_t __nv_cvt_float2_to_fp6x2(const float2 x, const
__nv_fp6_interpretation_t
fp6_interpretation, const
enum cudaRoundMode
rounding)
Converts input vector of two single precision numbers packed in float2 x into a vector of
two values of fp6 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
Converts input vector x to a vector of two fp6 values of the kind specified by
fp6_interpretation parameter, using rounding mode specified by rounding parame-
ter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result
in positive MAXNORM.
Returns
▶ The __nv_fp6x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp6_storage_t __nv_cvt_float_to_fp6(const float x, const
__nv_fp6_interpretation_t
fp6_interpretation, const enum
cudaRoundMode rounding)
Converts input single precision x to fp6 type of the requested kind using specified rounding
mode and saturating the out-of-range values.
Converts input x to fp6 type of the kind specified by fp6_interpretation parameter, us-
ing rounding mode specified by rounding parameter. Large out-of-range values saturate to
MAXNORM of the same sign. NaN input values result in positive MAXNORM.
Returns
▶ The __nv_fp6_storage_t value holds the result of conversion.
__host__ __device__ __half_raw __nv_cvt_fp6_to_halfraw(const __nv_fp6_storage_t x, const
__nv_fp6_interpretation_t
fp6_interpretation)
Converts input fp6 x of the specified kind to half precision.
Converts input x of fp6 type of the kind specified by fp6_interpretation parameter to half
precision.
Returns
▶ The __half_raw value holds the result of conversion.
__host__ __device__ __half2_raw __nv_cvt_fp6x2_to_halfraw2(const __nv_fp6x2_storage_t x,
const __nv_fp6_interpretation_t
fp6_interpretation)
Converts input vector of two fp6 values of the specified kind to a vector of two half precision
values packed in __half2_raw structure.
Converts input vector x of fp6 type of the kind specified by fp6_interpretation parameter
to a vector of two half precision values and returns as __half2_raw structure.
2.7. FP6 Conversion and Data Movement
21
Returns
▶ The __half2_raw value holds the result of conversion.
__host__ __device__ __nv_fp6x2_storage_t __nv_cvt_halfraw2_to_fp6x2(const __half2_raw x,
const
__nv_fp6_interpretation_t
fp6_interpretation,
const enum
cudaRoundMode
rounding)
Converts input vector of two half precision numbers packed in __half2_raw x into a vector of
two values of fp6 type of the requested kind using specified rounding mode and saturating the
out-of-range values.
Converts input vector x to a vector of two fp6 values of the kind specified by
fp6_interpretation parameter, using rounding mode specified by rounding parame-
ter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result
in positive MAXNORM.
Returns
▶ The __nv_fp6x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp6_storage_t __nv_cvt_halfraw_to_fp6(const __half_raw x, const
__nv_fp6_interpretation_t
fp6_interpretation, const
enum cudaRoundMode
rounding)
Converts input half precision x to fp6 type of the requested kind using specified rounding mode
and saturating the out-of-range values.
Converts input x to fp6 type of the kind specified by fp6_interpretation parameter, us-
ing rounding mode specified by rounding parameter. Large out-of-range values saturate to
MAXNORM of the same sign. NaN input values result in positive MAXNORM.
Returns
▶ The __nv_fp6_storage_t value holds the result of conversion.
2.7.3. Typedefs
typedef __nv_fp8_storage_t __nv_fp6_storage_t
8-bit unsigned integer type abstraction used for fp6 floating-point numbers storage.
typedef __nv_fp8x2_storage_t __nv_fp6x2_storage_t
16-bit unsigned integer type abstraction used for storage of pairs of fp6 floating-point num-
bers.
typedef __nv_fp8x4_storage_t __nv_fp6x4_storage_t
22
Chapter 2. FP6 Intrinsics
32-bit unsigned integer type abstraction used for storage of tetrads of fp6 floating-point
numbers.
Groups
C++ struct for handling fp6 data type of e2m3 kind.
C++ struct for handling fp6 data type of e3m2 kind.
C++ struct for handling vector type of four fp6 values of e2m3 kind.
C++ struct for handling vector type of four fp6 values of e3m2 kind.
C++ struct for handling vector type of two fp6 values of e2m3 kind.
C++ struct for handling vector type of two fp6 values of e3m2 kind.
FP6 Conversion and Data Movement
To use these functions, include the header file cuda_fp6.h in your program.
2.7. FP6 Conversion and Data Movement
23
24
Chapter 2. FP6 Intrinsics
Chapter 3. FP8 Intrinsics
This section describes fp8 intrinsic functions.
To use these functions, include the header file cuda_fp8.h in your program. The following macros are
available to help users selectively enable/disable various definitions present in the header file:
▶ __CUDA_NO_FP8_CONVERSIONS__ - If defined, this macro will prevent any use of the C++ type
conversions (converting constructors and conversion operators) defined in the header.
▶ __CUDA_NO_FP8_CONVERSION_OPERATORS__ - If defined, this macro will prevent any use of the
C++ conversion operators from fp8 to other types.
3.1. C++ struct for handling fp8 data type of
e4m3 kind.
Structs
__nv_fp8_e4m3
__nv_fp8_e4m3 datatype
3.2. C++ struct for handling fp8 data type of
e5m2 kind.
Structs
__nv_fp8_e5m2
__nv_fp8_e5m2 datatype
25
3.3. C++ struct for handling vector type of four
fp8 values of e4m3 kind.
Structs
__nv_fp8x4_e4m3
__nv_fp8x4_e4m3 datatype
3.4. C++ struct for handling vector type of four
fp8 values of e5m2 kind.
Structs
__nv_fp8x4_e5m2
__nv_fp8x4_e5m2 datatype
3.5. C++ struct for handling vector type of four
scale factors of e8m0 kind.
Structs
__nv_fp8x4_e8m0
__nv_fp8x4_e8m0 datatype
3.6. C++ struct for handling vector type of two
fp8 values of e4m3 kind.
Structs
__nv_fp8x2_e4m3
__nv_fp8x2_e4m3 datatype
26
Chapter 3. FP8 Intrinsics
3.7. C++ struct for handling vector type of two
fp8 values of e5m2 kind.
Structs
__nv_fp8x2_e5m2
__nv_fp8x2_e5m2 datatype
3.8. C++ struct for handling vector type of two
scale factors of e8m0 kind.
Structs
__nv_fp8x2_e8m0
__nv_fp8x2_e8m0 datatype
3.9. FP8 Conversion and Data Movement
To use these functions, include the header file cuda_fp8.h in your program.
Enumerations
__nv_fp8_interpretation_t
Enumerates the possible interpretations of the 8-bit values when referring to them as fp8 types.
__nv_saturation_t
Enumerates the modes applicable when performing a narrowing conversion to fp8 destination
types.
Functions
__host__
__device__
__nv_fp8x2_storage_t
__nv_cvt_bfloat162raw_to_e8m0x2(const
__nv_bfloat162_raw x, const
__nv_saturation_t saturate, const enum cudaRoundMode
rounding)
Converts a pair of bfloat16 values into a pair of scaling factors of e8m0 kind.
__host__
__device__
__nv_fp8x2_storage_t
__nv_cvt_bfloat16raw2_to_fp8x2(const
__nv_bfloat162_raw x, const __nv_saturation_t saturate, const __nv_fp8_interpretation_t
fp8_interpretation)
Converts input vector of two nv_bfloat16 precision numbers packed in __nv_bfloat162_raw
x into a vector of two values of fp8 type of the requested kind using round-to-nearest-even
rounding and requested saturation mode.
3.7. C++ struct for handling vector type of two fp8 values of e5m2 kind.
27
__host__ __device__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_e8m0(const __nv_bfloat16_raw
x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding)
Converts input bfloat16 input into a scaling factor of e8m0 kind.
__host__ __device__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_fp8(const __nv_bfloat16_raw x,
const __nv_saturation_t saturate, const __nv_fp8_interpretation_t fp8_interpretation)
Converts input nv_bfloat16 precision x to fp8 type of the requested kind using round-to-
nearest-even rounding and requested saturation mode.
__host__ __device__ __nv_fp8x2_storage_t __nv_cvt_double2_to_e8m0x2(const double2 x, const
__nv_saturation_t saturate, const enum cudaRoundMode rounding)
Converts a pair of double values into a pair of scaling factors of e8m0 kind.
__host__ __device__ __nv_fp8x2_storage_t __nv_cvt_double2_to_fp8x2(const double2 x, const
__nv_saturation_t saturate, const __nv_fp8_interpretation_t fp8_interpretation)
Converts input vector of two double precision numbers packed in double2 x into a vector of two
values of fp8 type of the requested kind using round-to-nearest-even rounding and requested
saturation mode.
__host__
__device__
__nv_fp8_storage_t
__nv_cvt_double_to_e8m0(const double x, const
__nv_saturation_t saturate, const enum cudaRoundMode rounding)
Converts input double value into a scaling factor of e8m0 kind.
__host__
__device__
__nv_fp8_storage_t
__nv_cvt_double_to_fp8(const double x,
const
__nv_saturation_t saturate, const __nv_fp8_interpretation_t fp8_interpretation)
Converts input double precision x to fp8 type of the requested kind using round-to-nearest-
even rounding and requested saturation mode.
__host__ __device__ __nv_bfloat16_raw __nv_cvt_e8m0_to_bf16raw(const __nv_fp8_storage_t x)
Converts input scaling factor value of e8m0 kind into bfloat16 .
__host__
__device__
__nv_bfloat162_raw
__nv_cvt_e8m0x2_to_bf162raw(const
__nv_fp8x2_storage_t x)
Converts input pair of scaling factors of e8m0 kind into a pair of bfloat16 values.
__host__
__device__
__nv_fp8x2_storage_t
__nv_cvt_float2_to_e8m0x2(const float2 x, const
__nv_saturation_t saturate, const enum cudaRoundMode rounding)
Converts a pair of float values into a pair of scaling factors of e8m0 kind.
__host__
__device__
__nv_fp8x2_storage_t
__nv_cvt_float2_to_fp8x2(const float2 x, const
__nv_saturation_t saturate, const __nv_fp8_interpretation_t fp8_interpretation)
Converts input vector of two single precision numbers packed in float2 x into a vector of two
values of fp8 type of the requested kind using round-to-nearest-even rounding and requested
saturation mode.
__host__
__device__
__nv_fp8_storage_t
__nv_cvt_float_to_e8m0(const
float
x,
const
__nv_saturation_t saturate, const enum cudaRoundMode rounding)
Converts input float value into a scaling factor of e8m0 kind.
__host__
__device__
__nv_fp8_storage_t
__nv_cvt_float_to_fp8(const
float
x,
const
__nv_saturation_t saturate, const __nv_fp8_interpretation_t fp8_interpretation)
Converts input single precision x to fp8 type of the requested kind using round-to-nearest-
even rounding and requested saturation mode.
__host__
__device__
__half_raw
__nv_cvt_fp8_to_halfraw(const
__nv_fp8_storage_t x, const
__nv_fp8_interpretation_t fp8_interpretation)
Converts input fp8 x of the specified kind to half precision.
__host__ __device__ __half2_raw __nv_cvt_fp8x2_to_halfraw2(const __nv_fp8x2_storage_t x, const
__nv_fp8_interpretation_t fp8_interpretation)
28
Chapter 3. FP8 Intrinsics
Converts input vector of two fp8 values of the specified kind to a vector of two half precision
values packed in __half2_raw structure.
__host__ __device__ __nv_fp8x2_storage_t __nv_cvt_halfraw2_to_fp8x2(const __half2_raw x, const
__nv_saturation_t saturate, const __nv_fp8_interpretation_t fp8_interpretation)
Converts input vector of two half precision numbers packed in __half2_raw x into a vector
of two values of fp8 type of the requested kind using round-to-nearest-even rounding and re-
quested saturation mode.
__host__
__device__
__nv_fp8_storage_t
__nv_cvt_halfraw_to_fp8(const
__half_raw x, const
__nv_saturation_t saturate, const __nv_fp8_interpretation_t fp8_interpretation)
Converts input half precision x to fp8 type of the requested kind using round-to-nearest-even
rounding and requested saturation mode.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const int val)
Constructor from int data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const long int val)
Constructor from long int data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const long long int val)
Constructor from long long int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__nv_fp8_e4m3::__nv_fp8_e4m3()=default
Constructor by default.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const unsigned short int val)
Constructor from unsigned short int data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const unsigned long int val)
Constructor from unsigned long int data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const short int val)
Constructor from short int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp8_e4m3::__nv_fp8_e4m3(const unsigned int val)
Constructor from unsigned int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
3.9. FP8 Conversion and Data Movement
29
__host__ __device__ __nv_fp8_e4m3::operator __half() const
Conversion operator to __half data type.
__host__ __device__ __nv_fp8_e4m3::operator __nv_bfloat16() const
Conversion operator to __nv_bfloat16 data type.
__host__ __device__ __nv_fp8_e4m3::operator bool() const
Conversion operator to bool data type.
__host__ __device__ __nv_fp8_e4m3::operator char() const
Conversion operator to an implementation defined char data type.
__host__ __device__ __nv_fp8_e4m3::operator double() const
Conversion operator to double data type.
__host__ __device__ __nv_fp8_e4m3::operator float() const
Conversion operator to float data type.
__host__ __device__ __nv_fp8_e4m3::operator int() const
Conversion operator to int data type.
__host__ __device__ __nv_fp8_e4m3::operator long int() const
Conversion operator to long int data type.
__host__ __device__ __nv_fp8_e4m3::operator long long int() const
Conversion operator to long long int data type.
__host__ __device__ __nv_fp8_e4m3::operator short int() const
Conversion operator to short int data type.
__host__ __device__ __nv_fp8_e4m3::operator signed char() const
Conversion operator to signed char data type.
__host__ __device__ __nv_fp8_e4m3::operator unsigned char() const
Conversion operator to unsigned char data type.
__host__ __device__ __nv_fp8_e4m3::operator unsigned int() const
Conversion operator to unsigned int data type.
__host__ __device__ __nv_fp8_e4m3::operator unsigned long int() const
Conversion operator to unsigned long int data type.
__host__ __device__ __nv_fp8_e4m3::operator unsigned long long int() const
Conversion operator to unsigned long long int data type.
__host__ __device__ __nv_fp8_e4m3::operator unsigned short int() const
Conversion operator to unsigned short int data type.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const long long int val)
Constructor from long long int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const unsigned int val)
Constructor from unsigned int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior for out-of-range values.
30
Chapter 3. FP8 Intrinsics
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const unsigned short int val)
Constructor from unsigned short int data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__nv_fp8_e5m2::__nv_fp8_e5m2()=default
Constructor by default.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const int val)
Constructor from int data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const long int val)
Constructor from long int data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const unsigned long int val)
Constructor from unsigned long int data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const short int val)
Constructor from short int data type.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ __nv_fp8_e5m2::__nv_fp8_e5m2(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ __nv_fp8_e5m2::operator __half() const
Conversion operator to __half data type.
__host__ __device__ __nv_fp8_e5m2::operator __nv_bfloat16() const
Conversion operator to __nv_bfloat16 data type.
__host__ __device__ __nv_fp8_e5m2::operator bool() const
Conversion operator to bool data type.
__host__ __device__ __nv_fp8_e5m2::operator char() const
Conversion operator to an implementation defined char data type.
__host__ __device__ __nv_fp8_e5m2::operator double() const
Conversion operator to double data type.
__host__ __device__ __nv_fp8_e5m2::operator float() const
Conversion operator to float data type.
__host__ __device__ __nv_fp8_e5m2::operator int() const
Conversion operator to int data type.
__host__ __device__ __nv_fp8_e5m2::operator long int() const
Conversion operator to long int data type.
__host__ __device__ __nv_fp8_e5m2::operator long long int() const
Conversion operator to long long int data type.
__host__ __device__ __nv_fp8_e5m2::operator short int() const
Conversion operator to short int data type.
3.9. FP8 Conversion and Data Movement
31
__host__ __device__ __nv_fp8_e5m2::operator signed char() const
Conversion operator to signed char data type.
__host__ __device__ __nv_fp8_e5m2::operator unsigned char() const
Conversion operator to unsigned char data type.
__host__ __device__ __nv_fp8_e5m2::operator unsigned int() const
Conversion operator to unsigned int data type.
__host__ __device__ __nv_fp8_e5m2::operator unsigned long int() const
Conversion operator to unsigned long int data type.
__host__ __device__ __nv_fp8_e5m2::operator unsigned long long int() const
Conversion operator to unsigned long long int data type.
__host__ __device__ __nv_fp8_e5m2::operator unsigned short int() const
Conversion operator to unsigned short int data type.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const long int val)
Constructor from long int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const int val)
Constructor from int data type, relies on cudaRoundPosInf rounding.
__nv_fp8_e8m0::__nv_fp8_e8m0()=default
Constructor by default.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const unsigned int val)
Constructor from unsigned int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior behavior for large input
values and cudaRoundPosInf for rounding.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for large input values
and cudaRoundPosInf for rounding.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for large input values
and cudaRoundPosInf for rounding.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for large input
values and cudaRoundPosInf for rounding.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const unsigned long int val)
Constructor from unsigned long int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const unsigned short int val)
Constructor from unsigned short int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const long long int val)
Constructor from long long int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ __nv_fp8_e8m0::__nv_fp8_e8m0(const short int val)
Constructor from short int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ __nv_fp8_e8m0::operator __half() const
Conversion operator to __half data type.
32
Chapter 3. FP8 Intrinsics
__host__ __device__ __nv_fp8_e8m0::operator __nv_bfloat16() const
Conversion operator to __nv_bfloat16 data type.
__host__ __device__ __nv_fp8_e8m0::operator bool() const
Conversion operator to bool data type.
__host__ __device__ __nv_fp8_e8m0::operator char() const
Conversion operator to an implementation defined char data type.
__host__ __device__ __nv_fp8_e8m0::operator double() const
Conversion operator to double data type.
__host__ __device__ __nv_fp8_e8m0::operator float() const
Conversion operator to float data type.
__host__ __device__ __nv_fp8_e8m0::operator int() const
Conversion operator to int data type.
__host__ __device__ __nv_fp8_e8m0::operator long int() const
Conversion operator to long int data type.
__host__ __device__ __nv_fp8_e8m0::operator long long int() const
Conversion operator to long long int data type.
__host__ __device__ __nv_fp8_e8m0::operator short int() const
Conversion operator to short int data type.
__host__ __device__ __nv_fp8_e8m0::operator signed char() const
Conversion operator to signed char data type.
__host__ __device__ __nv_fp8_e8m0::operator unsigned char() const
Conversion operator to unsigned char data type.
__host__ __device__ __nv_fp8_e8m0::operator unsigned int() const
Conversion operator to unsigned int data type.
__host__ __device__ __nv_fp8_e8m0::operator unsigned long int() const
Conversion operator to unsigned long int data type.
__host__ __device__ __nv_fp8_e8m0::operator unsigned long long int() const
Conversion operator to unsigned long long int data type.
__host__ __device__ __nv_fp8_e8m0::operator unsigned short int() const
Conversion operator to unsigned short int data type.
__nv_fp8x2_e4m3::__nv_fp8x2_e4m3()=default
Constructor by default.
__host__ __device__ __nv_fp8x2_e4m3::__nv_fp8x2_e4m3(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ __nv_fp8x2_e4m3::__nv_fp8x2_e4m3(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp8x2_e4m3::__nv_fp8x2_e4m3(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp8x2_e4m3::__nv_fp8x2_e4m3(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range values.
3.9. FP8 Conversion and Data Movement
33
__host__ __device__ __nv_fp8x2_e4m3::operator __half2() const
Conversion operator to __half2 data type.
__host__ __device__ __nv_fp8x2_e4m3::operator float2() const
Conversion operator to float2 data type.
__host__ __device__ __nv_fp8x2_e5m2::__nv_fp8x2_e5m2(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ __nv_fp8x2_e5m2::__nv_fp8x2_e5m2(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp8x2_e5m2::__nv_fp8x2_e5m2(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__nv_fp8x2_e5m2::__nv_fp8x2_e5m2()=default
Constructor by default.
__host__ __device__ __nv_fp8x2_e5m2::__nv_fp8x2_e5m2(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp8x2_e5m2::operator __half2() const
Conversion operator to __half2 data type.
__host__ __device__ __nv_fp8x2_e5m2::operator float2() const
Conversion operator to float2 data type.
__host__ __device__ __nv_fp8x2_e8m0::__nv_fp8x2_e8m0(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ __nv_fp8x2_e8m0::__nv_fp8x2_e8m0(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range values.
__host__ __device__ __nv_fp8x2_e8m0::__nv_fp8x2_e8m0(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ __nv_fp8x2_e8m0::__nv_fp8x2_e8m0(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__nv_fp8x2_e8m0::__nv_fp8x2_e8m0()=default
Constructor by default.
__host__ __device__ __nv_fp8x2_e8m0::operator __half2() const
Conversion operator to __half2 data type.
__host__ __device__ __nv_fp8x2_e8m0::operator __nv_bfloat162() const
Conversion operator to __nv_bfloat162 data type.
__host__ __device__ __nv_fp8x2_e8m0::operator float2() const
Conversion operator to float2 data type.
__host__
__device__
__nv_fp8x4_e4m3::__nv_fp8x4_e4m3(const
__nv_bfloat162 flo,
const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
34
Chapter 3. FP8 Intrinsics
__host__ __device__ __nv_fp8x4_e4m3::__nv_fp8x4_e4m3(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__nv_fp8x4_e4m3::__nv_fp8x4_e4m3()=default
Constructor by default.
__host__ __device__ __nv_fp8x4_e4m3::__nv_fp8x4_e4m3(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp8x4_e4m3::__nv_fp8x4_e4m3(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ __nv_fp8x4_e4m3::operator float4() const
Conversion operator to float4 vector data type.
__nv_fp8x4_e5m2::__nv_fp8x4_e5m2()=default
Constructor by default.
__host__ __device__ __nv_fp8x4_e5m2::__nv_fp8x4_e5m2(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ __nv_fp8x4_e5m2::__nv_fp8x4_e5m2(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__
__device__
__nv_fp8x4_e5m2::__nv_fp8x4_e5m2(const
__nv_bfloat162 flo,
const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
__host__ __device__ __nv_fp8x4_e5m2::__nv_fp8x4_e5m2(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ __nv_fp8x4_e5m2::operator float4() const
Conversion operator to float4 vector data type.
__host__ __device__ __nv_fp8x4_e8m0::__nv_fp8x4_e8m0(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ __nv_fp8x4_e8m0::__nv_fp8x4_e8m0(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ __nv_fp8x4_e8m0::__nv_fp8x4_e8m0(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__
__device__
__nv_fp8x4_e8m0::__nv_fp8x4_e8m0(const
__nv_bfloat162 flo,
const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
__nv_fp8x4_e8m0::__nv_fp8x4_e8m0()=default
Constructor by default.
3.9. FP8 Conversion and Data Movement
35
__host__ __device__ __nv_fp8x4_e8m0::operator float4() const
Conversion operator to float4 vector data type.
Typedefs
__nv_fp8_storage_t
8-bit unsigned integer type abstraction used for fp8 floating-point numbers storage.
__nv_fp8x2_storage_t
16-bit unsigned integer type abstraction used for storage of pairs of fp8 floating-point num-
bers.
__nv_fp8x4_storage_t
32-bit unsigned integer type abstraction used for storage of tetrads of fp8 floating-point
numbers.
3.9.1. Enumerations
enum __nv_fp8_interpretation_t
Enumerates the possible interpretations of the 8-bit values when referring to them as fp8 types.
Values:
enumerator __NV_E4M3
Stands for fp8 numbers of e4m3 kind.
enumerator __NV_E5M2
Stands for fp8 numbers of e5m2 kind.
enum __nv_saturation_t
Enumerates the modes applicable when performing a narrowing conversion to fp8 destination
types.
Values:
enumerator __NV_NOSAT
Means no saturation to finite is performed when conversion results in rounding values out-
side the range of destination type.
NOTE: for fp8 type of e4m3 kind, the results that are larger than the maximum representable
finite number of the target format become NaN.
enumerator __NV_SATFINITE
Means input larger than the maximum representable finite number MAXNORM of the target
format round to the MAXNORM of the same sign as input.
36
Chapter 3. FP8 Intrinsics
3.9.2. Functions
__host__ __device__ __nv_fp8x2_storage_t __nv_cvt_bfloat162raw_to_e8m0x2(const
__nv_bfloat162_raw
x, const
__nv_saturation_t
saturate, const
enum cud-
aRoundMode
rounding)
Converts a pair of bfloat16 values into a pair of scaling factors of e8m0 kind.
See also:
__nv_cvt_bfloat16raw_to_e8m0() for details of conversion.
Returns
▶ The __nv_fp8x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp8x2_storage_t __nv_cvt_bfloat16raw2_to_fp8x2(const
__nv_bfloat162_raw
x, const
__nv_saturation_t
saturate, const
__nv_fp8_interpretation_t
fp8_interpretation)
Converts input vector of two nv_bfloat16 precision numbers packed in __nv_bfloat162_raw
x into a vector of two values of fp8 type of the requested kind using round-to-nearest-even
rounding and requested saturation mode.
Converts input vector x to a vector of two fp8 values of the kind specified by
fp8_interpretation parameter, using round-to-nearest-even rounding and saturation
mode specified by saturate parameter.
Returns
▶ The __nv_fp8x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_e8m0(const
__nv_bfloat16_raw x,
const
__nv_saturation_t
saturate, const enum
cudaRoundMode
rounding)
Converts input bfloat16 input into a scaling factor of e8m0 kind.
Input number’s absolute value is rounded to the closest power of two in the direction specified
via rounding parameter. Rounded results that are smaller than the smallest representable tar-
get format number 2^-127 are then clipped to 2^-127. Results that are larger than the largest
representable target format number 2^127 are either clipped to 2^127 if saturate equals to
3.9. FP8 Conversion and Data Movement
37
__NV_SATFINITE, or convert to NaN otherwise. NaN inputs convert into NaN output, encoded as
0xFF in the target format.
Returns
▶ The __nv_fp8_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_fp8(const
__nv_bfloat16_raw x,
const __nv_saturation_t
saturate, const
__nv_fp8_interpretation_t
fp8_interpretation)
Converts input nv_bfloat16 precision x to fp8 type of the requested kind using round-to-
nearest-even rounding and requested saturation mode.
Converts input x to fp8 type of the kind specified by fp8_interpretation parameter, using
round-to-nearest-even rounding and saturation mode specified by saturate parameter.
Returns
▶ The __nv_fp8_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp8x2_storage_t __nv_cvt_double2_to_e8m0x2(const double2 x, const
__nv_saturation_t
saturate, const enum
cudaRoundMode
rounding)
Converts a pair of double values into a pair of scaling factors of e8m0 kind.
See also:
__nv_cvt_bfloat16raw_to_e8m0() for details of conversion.
Returns
▶ The __nv_fp8x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp8x2_storage_t __nv_cvt_double2_to_fp8x2(const double2 x, const
__nv_saturation_t
saturate, const
__nv_fp8_interpretation_t
fp8_interpretation)
Converts input vector of two double precision numbers packed in double2 x into a vector of two
values of fp8 type of the requested kind using round-to-nearest-even rounding and requested
saturation mode.
Converts input vector x to a vector of two fp8 values of the kind specified by
fp8_interpretation parameter, using round-to-nearest-even rounding and saturation
mode specified by saturate parameter.
Returns
▶ The __nv_fp8x2_storage_t value holds the result of conversion.
38
Chapter 3. FP8 Intrinsics
__host__ __device__ __nv_fp8_storage_t __nv_cvt_double_to_e8m0(const double x, const
__nv_saturation_t saturate,
const enum cudaRoundMode
rounding)
Converts input double value into a scaling factor of e8m0 kind.
See also:
__nv_cvt_bfloat16raw_to_e8m0() for details of conversion.
Returns
▶ The __nv_fp8_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp8_storage_t __nv_cvt_double_to_fp8(const double x, const
__nv_saturation_t saturate,
const
__nv_fp8_interpretation_t
fp8_interpretation)
Converts input double precision x to fp8 type of the requested kind using round-to-nearest-
even rounding and requested saturation mode.
Converts input x to fp8 type of the kind specified by fp8_interpretation parameter, using
round-to-nearest-even rounding and saturation mode specified by saturate parameter.
Returns
▶ The __nv_fp8_storage_t value holds the result of conversion.
__host__ __device__ __nv_bfloat16_raw __nv_cvt_e8m0_to_bf16raw(const __nv_fp8_storage_t x)
Converts input scaling factor value of e8m0 kind into bfloat16.
Input scales are exact powers of two or a NaN value, also representable in the target format.
Returns
▶ The __nv_bfloat16_raw value holds the result of conversion.
__host__ __device__ __nv_bfloat162_raw __nv_cvt_e8m0x2_to_bf162raw(const
__nv_fp8x2_storage_t
x)
Converts input pair of scaling factors of e8m0 kind into a pair of bfloat16 values.
Returns
▶ The __nv_bfloat162_raw value holds the result of conversion.
__host__ __device__ __nv_fp8x2_storage_t __nv_cvt_float2_to_e8m0x2(const float2 x, const
__nv_saturation_t
saturate, const enum
cudaRoundMode
rounding)
Converts a pair of float values into a pair of scaling factors of e8m0 kind.
3.9. FP8 Conversion and Data Movement
39
See also:
__nv_cvt_bfloat16raw_to_e8m0() for details of conversion.
Returns
▶ The __nv_fp8x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp8x2_storage_t __nv_cvt_float2_to_fp8x2(const float2 x, const
__nv_saturation_t
saturate, const
__nv_fp8_interpretation_t
fp8_interpretation)
Converts input vector of two single precision numbers packed in float2 x into a vector of two
values of fp8 type of the requested kind using round-to-nearest-even rounding and requested
saturation mode.
Converts input vector x to a vector of two fp8 values of the kind specified by
fp8_interpretation parameter, using round-to-nearest-even rounding and saturation
mode specified by saturate parameter.
Returns
▶ The __nv_fp8x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp8_storage_t __nv_cvt_float_to_e8m0(const float x, const
__nv_saturation_t saturate,
const enum cudaRoundMode
rounding)
Converts input float value into a scaling factor of e8m0 kind.
See also:
__nv_cvt_bfloat16raw_to_e8m0() for details of conversion.
Returns
▶ The __nv_fp8_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp8_storage_t __nv_cvt_float_to_fp8(const float x, const
__nv_saturation_t saturate,
const __nv_fp8_interpretation_t
fp8_interpretation)
Converts input single precision x to fp8 type of the requested kind using round-to-nearest-
even rounding and requested saturation mode.
Converts input x to fp8 type of the kind specified by fp8_interpretation parameter, using
round-to-nearest-even rounding and saturation mode specified by saturate parameter.
Returns
▶ The __nv_fp8_storage_t value holds the result of conversion.
40
Chapter 3. FP8 Intrinsics
__host__ __device__ __half_raw __nv_cvt_fp8_to_halfraw(const __nv_fp8_storage_t x, const
__nv_fp8_interpretation_t
fp8_interpretation)
Converts input fp8 x of the specified kind to half precision.
Converts input x of fp8 type of the kind specified by fp8_interpretation parameter to half
precision.
Returns
▶ The __half_raw value holds the result of conversion.
__host__ __device__ __half2_raw __nv_cvt_fp8x2_to_halfraw2(const __nv_fp8x2_storage_t x,
const __nv_fp8_interpretation_t
fp8_interpretation)
Converts input vector of two fp8 values of the specified kind to a vector of two half precision
values packed in __half2_raw structure.
Converts input vector x of fp8 type of the kind specified by fp8_interpretation parameter
to a vector of two half precision values and returns as __half2_raw structure.
Returns
▶ The __half2_raw value holds the result of conversion.
__host__ __device__ __nv_fp8x2_storage_t __nv_cvt_halfraw2_to_fp8x2(const __half2_raw x,
const
__nv_saturation_t
saturate, const
__nv_fp8_interpretation_t
fp8_interpretation)
Converts input vector of two half precision numbers packed in __half2_raw x into a vector
of two values of fp8 type of the requested kind using round-to-nearest-even rounding and re-
quested saturation mode.
Converts input vector x to a vector of two fp8 values of the kind specified by
fp8_interpretation parameter, using round-to-nearest-even rounding and saturation
mode specified by saturate parameter.
Returns
▶ The __nv_fp8x2_storage_t value holds the result of conversion.
__host__ __device__ __nv_fp8_storage_t __nv_cvt_halfraw_to_fp8(const __half_raw x, const
__nv_saturation_t saturate,
const
__nv_fp8_interpretation_t
fp8_interpretation)
Converts input half precision x to fp8 type of the requested kind using round-to-nearest-even
rounding and requested saturation mode.
Converts input x to fp8 type of the kind specified by fp8_interpretation parameter, using
round-to-nearest-even rounding and saturation mode specified by saturate parameter.
Returns
3.9. FP8 Conversion and Data Movement
41
▶ The __nv_fp8_storage_t value holds the result of conversion.
3.9.3. Typedefs
typedef unsigned char __nv_fp8_storage_t
8-bit unsigned integer type abstraction used for fp8 floating-point numbers storage.
typedef unsigned short int __nv_fp8x2_storage_t
16-bit unsigned integer type abstraction used for storage of pairs of fp8 floating-point num-
bers.
typedef unsigned int __nv_fp8x4_storage_t
32-bit unsigned integer type abstraction used for storage of tetrads of fp8 floating-point
numbers.
Groups
C++ struct for handling fp8 data type of e4m3 kind.
C++ struct for handling fp8 data type of e5m2 kind.
C++ struct for handling vector type of four fp8 values of e4m3 kind.
C++ struct for handling vector type of four fp8 values of e5m2 kind.
C++ struct for handling vector type of four scale factors of e8m0 kind.
C++ struct for handling vector type of two fp8 values of e4m3 kind.
C++ struct for handling vector type of two fp8 values of e5m2 kind.
C++ struct for handling vector type of two scale factors of e8m0 kind.
FP8 Conversion and Data Movement
To use these functions, include the header file cuda_fp8.h in your program.
42
Chapter 3. FP8 Intrinsics
Chapter 4. Half Precision Intrinsics
This section describes half precision intrinsic functions.
To use these functions, include the header file cuda_fp16.h in your program. All of the functions
defined here are available in device code. Some of the functions are also available to host compilers,
please refer to respective functions’ documentation for details.
NOTE: Aggressive floating-point optimizations performed by host or device compilers may affect nu-
meric behavior of the functions implemented in this header.
The following macros are available to help users selectively enable/disable various definitions present
in the header file:
▶ CUDA_NO_HALF - If defined, this macro will prevent the definition of additional type aliases in the
global namespace, helping to avoid potential conflicts with symbols defined in the user program.
▶ __CUDA_NO_HALF_CONVERSIONS__ - If defined, this macro will prevent the use of the C++ type
conversions (converting constructors and conversion operators) that are common for built-in
floating-point types, but may be undesirable for half which is essentially a user-defined type.
▶ __CUDA_NO_HALF_OPERATORS__ and __CUDA_NO_HALF2_OPERATORS__ - If defined, these
macros will prevent the inadvertent use of usual arithmetic and comparison operators. This
enforces the storage-only type semantics and prevents C++ style computations on half and
half2 types.
4.1. Half Arithmetic Constants
To use these constants, include the header file cuda_fp16.h in your program.
Macros
CUDART_INF_FP16
Defines floating-point positive infinity value for the half data type.
CUDART_MAX_NORMAL_FP16
Defines a maximum representable value for the half data type.
CUDART_MIN_DENORM_FP16
Defines a minimum representable (denormalized) value for the half data type.
CUDART_NAN_FP16
Defines canonical NaN value for the half data type.
43
CUDART_NEG_ZERO_FP16
Defines a negative zero value for the half data type.
CUDART_ONE_FP16
Defines a value of 1.0 for the half data type.
CUDART_ZERO_FP16
Defines a positive zero value for the half data type.
4.1.1. Macros
CUDART_INF_FP16 __ushort_as_half((unsigned short)0x7C00U)
Defines floating-point positive infinity value for the half data type.
CUDART_MAX_NORMAL_FP16 __ushort_as_half((unsigned short)0x7BFFU)
Defines a maximum representable value for the half data type.
CUDART_MIN_DENORM_FP16 __ushort_as_half((unsigned short)0x0001U)
Defines a minimum representable (denormalized) value for the half data type.
CUDART_NAN_FP16 __ushort_as_half((unsigned short)0x7FFFU)
Defines canonical NaN value for the half data type.
CUDART_NEG_ZERO_FP16 __ushort_as_half((unsigned short)0x8000U)
Defines a negative zero value for the half data type.
CUDART_ONE_FP16 __ushort_as_half((unsigned short)0x3C00U)
Defines a value of 1.0 for the half data type.
CUDART_ZERO_FP16 __ushort_as_half((unsigned short)0x0000U)
Defines a positive zero value for the half data type.
4.2. Half Arithmetic Functions
To use these functions, include the header file cuda_fp16.h in your program.
44
Chapter 4. Half Precision Intrinsics
Functions
__host__ __device__ __half __habs(const __half a)
Calculates the absolute value of input half number and returns the result.
__host__ __device__ __half __hadd(const __half a, const __half b)
Performs half addition in round-to-nearest-even mode.
__host__ __device__ __half __hadd_rn(const __half a, const __half b)
Performs half addition in round-to-nearest-even mode.
__host__ __device__ __half __hadd_sat(const __half a, const __half b)
Performs half addition in round-to-nearest-even mode, with saturation to [0.0, 1.0].
__host__ __device__ __half __hdiv(const __half a, const __half b)
Performs half division in round-to-nearest-even mode.
__device__ __half __hfma(const __half a, const __half b, const __half c)
Performs half fused multiply-add in round-to-nearest-even mode.
__device__ __half __hfma_relu(const __half a, const __half b, const __half c)
Performs half fused multiply-add in round-to-nearest-even mode with relu saturation.
__device__ __half __hfma_sat(const __half a, const __half b, const __half c)
Performs half fused multiply-add in round-to-nearest-even mode, with saturation to [0.0, 1.0].
__host__ __device__ __half __hmul(const __half a, const __half b)
Performs half multiplication in round-to-nearest-even mode.
__host__ __device__ __half __hmul_rn(const __half a, const __half b)
Performs half multiplication in round-to-nearest-even mode.
__host__ __device__ __half __hmul_sat(const __half a, const __half b)
Performs half multiplication in round-to-nearest-even mode, with saturation to [0.0, 1.0].
__host__ __device__ __half __hneg(const __half a)
Negates input half number and returns the result.
__host__ __device__ __half __hsub(const __half a, const __half b)
Performs half subtraction in round-to-nearest-even mode.
__host__ __device__ __half __hsub_rn(const __half a, const __half b)
Performs half subtraction in round-to-nearest-even mode.
__host__ __device__ __half __hsub_sat(const __half a, const __half b)
Performs half subtraction in round-to-nearest-even mode, with saturation to [0.0, 1.0].
__device__ __half atomicAdd(__half *const address, const __half val)
Adds val to the value stored at address in global or shared memory, and writes this value back
to address .
__host__ __device__ __half operator*(const __half &lh, const __half &rh)
Performs half multiplication operation.
__host__ __device__ __half & operator*=(__half &lh, const __half &rh)
Performs half compound assignment with multiplication operation.
__host__ __device__ __half operator+(const __half &h)
Implements half unary plus operator, returns input value.
__host__ __device__ __half operator+(const __half &lh, const __half &rh)
Performs half addition operation.
4.2. Half Arithmetic Functions
45
__host__ __device__ __half & operator++(__half &h)
Performs half prefix increment operation.
__host__ __device__ __half operator++(__half &h, const int ignored)
Performs half postfix increment operation.
__host__ __device__ __half & operator+=(__half &lh, const __half &rh)
Performs half compound assignment with addition operation.
__host__ __device__ __half operator-(const __half &lh, const __half &rh)
Performs half subtraction operation.
__host__ __device__ __half operator-(const __half &h)
Implements half unary minus operator.
__host__ __device__ __half operator-(__half &h, const int ignored)
Performs half postfix decrement operation.
__host__ __device__ __half & operator-(__half &h)
Performs half prefix decrement operation.
__host__ __device__ __half & operator-=(__half &lh, const __half &rh)
Performs half compound assignment with subtraction operation.
__host__ __device__ __half operator/(const __half &lh, const __half &rh)
Performs half division operation.
__host__ __device__ __half & operator/=(__half &lh, const __half &rh)
Performs half compound assignment with division operation.
4.2.1. Functions
__host__ __device__ __half __habs(const __half a)
Calculates the absolute value of input half number and returns the result.
Calculates the absolute value of input half number and returns the result.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The absolute value of a.
▶ __habs (±0) returns +0.
▶ __habs (±∞) returns +∞.
▶ __habs(NaN) returns NaN.
__host__ __device__ __half __hadd(const __half a, const __half b)
Performs half addition in round-to-nearest-even mode.
Performs half addition of inputs a and b, in round-to-nearest-even mode.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
46
Chapter 4. Half Precision Intrinsics
Returns
half
▶ The sum of a and b.
__host__ __device__ __half __hadd_rn(const __half a, const __half b)
Performs half addition in round-to-nearest-even mode.
Performs half addition of inputs a and b, in round-to-nearest-even mode. Prevents floating-
point contractions of mul+add into fma.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
▶ The sum of a and b.
__host__ __device__ __half __hadd_sat(const __half a, const __half b)
Performs half addition in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs half add of inputs a and b, in round-to-nearest-even mode, and clamps the result to
range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
▶ The sum of a and b, with respect to saturation.
__host__ __device__ __half __hdiv(const __half a, const __half b)
Performs half division in round-to-nearest-even mode.
Divides half input a by input b in round-to-nearest-even mode.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
▶ The result of dividing a by b.
__device__ __half __hfma(const __half a, const __half b, const __half c)
Performs half fused multiply-add in round-to-nearest-even mode.
Performs half multiply on inputs a and b, then performs a half add of the result with c, rounding
the result once in round-to-nearest-even mode.
Parameters
4.2. Half Arithmetic Functions
47
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
▶ c - [in] - half. Is only being read.
Returns
half
▶ The result of fused multiply-add operation on a, b, and c.
__device__ __half __hfma_relu(const __half a, const __half b, const __half c)
Performs half fused multiply-add in round-to-nearest-even mode with relu saturation.
Performs half multiply on inputs a and b, then performs a half add of the result with c, rounding
the result once in round-to-nearest-even mode. Then negative result is clamped to 0. NaN result
is converted to canonical NaN.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
▶ c - [in] - half. Is only being read.
Returns
half
▶ The result of fused multiply-add operation on a, b, and c with relu saturation.
__device__ __half __hfma_sat(const __half a, const __half b, const __half c)
Performs half fused multiply-add in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs half multiply on inputs a and b, then performs a half add of the result with c, rounding
the result once in round-to-nearest-even mode, and clamps the result to range [0.0, 1.0]. NaN
results are flushed to +0.0.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
▶ c - [in] - half. Is only being read.
Returns
half
▶ The result of fused multiply-add operation on a, b, and c, with respect to satu-
ration.
__host__ __device__ __half __hmul(const __half a, const __half b)
Performs half multiplication in round-to-nearest-even mode.
Performs half multiplication of inputs a and b, in round-to-nearest-even mode.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
48
Chapter 4. Half Precision Intrinsics
Returns
half
▶ The result of multiplying a and b.
__host__ __device__ __half __hmul_rn(const __half a, const __half b)
Performs half multiplication in round-to-nearest-even mode.
Performs half multiplication of inputs a and b, in round-to-nearest-even mode. Prevents
floating-point contractions of mul+add or sub into fma.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
▶ The result of multiplying a and b.
__host__ __device__ __half __hmul_sat(const __half a, const __half b)
Performs half multiplication in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs half multiplication of inputs a and b, in round-to-nearest-even mode, and clamps the
result to range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
▶ The result of multiplying a and b, with respect to saturation.
__host__ __device__ __half __hneg(const __half a)
Negates input half number and returns the result.
Negates input half number and returns the result.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ Negated input a.
▶ __hneg (±0) returns ∓0.
▶ __hneg (±∞) returns ∓∞.
▶ __hneg(NaN) returns NaN.
__host__ __device__ __half __hsub(const __half a, const __half b)
Performs half subtraction in round-to-nearest-even mode.
Subtracts half input b from input a in round-to-nearest-even mode.
4.2. Half Arithmetic Functions
49
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
▶ The result of subtracting b from a.
__host__ __device__ __half __hsub_rn(const __half a, const __half b)
Performs half subtraction in round-to-nearest-even mode.
Subtracts half input b from input a in round-to-nearest-even mode. Prevents floating-point
contractions of mul+sub into fma.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
▶ The result of subtracting b from a.
__host__ __device__ __half __hsub_sat(const __half a, const __half b)
Performs half subtraction in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Subtracts half input b from input a in round-to-nearest-even mode, and clamps the result to
range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
▶ The result of subtraction of b from a, with respect to saturation.
__device__ __half atomicAdd(__half *const address, const __half val)
Adds val to the value stored at address in global or shared memory, and writes this value back
to address.
This operation is performed in one atomic operation.
The location of address must be in global or shared memory. This operation has undefined
behavior otherwise. This operation is only supported by devices of compute capability 7.x and
higher.
Note: For more details about this function, see the Atomic Functions section in the CUDA C++
Programming Guide.
50
Chapter 4. Half Precision Intrinsics
Parameters
▶ address - [in] - half*. An address in global or shared memory.
▶ val - [in] - half. The value to be added.
Returns
half
▶ The old value read from address.
__host__ __device__ __half operator*(const __half &lh, const __half &rh)
Performs half multiplication operation.
See also:
__hmul(__half, __half)
__host__ __device__ __half &operator*=(__half &lh, const __half &rh)
Performs half compound assignment with multiplication operation.
See also:
__hmul(__half, __half)
__host__ __device__ __half operator+(const __half &h)
Implements half unary plus operator, returns input value.
__host__ __device__ __half operator+(const __half &lh, const __half &rh)
Performs half addition operation.
See also:
__hadd(__half, __half)
__host__ __device__ __half &operator++(__half &h)
Performs half prefix increment operation.
See also:
__hadd(__half, __half)
__host__ __device__ __half operator++(__half &h, const int ignored)
Performs half postfix increment operation.
See also:
__hadd(__half, __half)
__host__ __device__ __half &operator+=(__half &lh, const __half &rh)
Performs half compound assignment with addition operation.
4.2. Half Arithmetic Functions
51
See also:
__hadd(__half, __half)
__host__ __device__ __half operator-(const __half &lh, const __half &rh)
Performs half subtraction operation.
See also:
__hsub(__half, __half)
__host__ __device__ __half operator-(const __half &h)
Implements half unary minus operator.
See also:
__hneg(__half)
__host__ __device__ __half operator--(__half &h, const int ignored)
Performs half postfix decrement operation.
See also:
__hsub(__half, __half)
__host__ __device__ __half &operator--(__half &h)
Performs half prefix decrement operation.
See also:
__hsub(__half, __half)
__host__ __device__ __half &operator-=(__half &lh, const __half &rh)
Performs half compound assignment with subtraction operation.
See also:
__hsub(__half, __half)
__host__ __device__ __half operator∕(const __half &lh, const __half &rh)
Performs half division operation.
See also:
__hdiv(__half, __half)
__host__ __device__ __half &operator∕=(__half &lh, const __half &rh)
Performs half compound assignment with division operation.
See also:
__hdiv(__half, __half)
52
Chapter 4. Half Precision Intrinsics
4.3. Half Comparison Functions
To use these functions, include the header file cuda_fp16.h in your program.
Functions
__host__ __device__ bool __heq(const __half a, const __half b)
Performs half if-equal comparison.
__host__ __device__ bool __hequ(const __half a, const __half b)
Performs half unordered if-equal comparison.
__host__ __device__ bool __hge(const __half a, const __half b)
Performs half greater-equal comparison.
__host__ __device__ bool __hgeu(const __half a, const __half b)
Performs half unordered greater-equal comparison.
__host__ __device__ bool __hgt(const __half a, const __half b)
Performs half greater-than comparison.
__host__ __device__ bool __hgtu(const __half a, const __half b)
Performs half unordered greater-than comparison.
__host__ __device__ int __hisinf(const __half a)
Checks if the input half number is infinite.
__host__ __device__ bool __hisnan(const __half a)
Determine whether half argument is a NaN.
__host__ __device__ bool __hle(const __half a, const __half b)
Performs half less-equal comparison.
__host__ __device__ bool __hleu(const __half a, const __half b)
Performs half unordered less-equal comparison.
__host__ __device__ bool __hlt(const __half a, const __half b)
Performs half less-than comparison.
__host__ __device__ bool __hltu(const __half a, const __half b)
Performs half unordered less-than comparison.
__host__ __device__ __half __hmax(const __half a, const __half b)
Calculates half maximum of two input values.
__host__ __device__ __half __hmax_nan(const __half a, const __half b)
Calculates half maximum of two input values, NaNs pass through.
__host__ __device__ __half __hmin(const __half a, const __half b)
Calculates half minimum of two input values.
__host__ __device__ __half __hmin_nan(const __half a, const __half b)
Calculates half minimum of two input values, NaNs pass through.
__host__ __device__ bool __hne(const __half a, const __half b)
Performs half not-equal comparison.
__host__ __device__ bool __hneu(const __half a, const __half b)
Performs half unordered not-equal comparison.
4.3. Half Comparison Functions
53
__host__ __device__ bool operator!=(const __half &lh, const __half &rh)
Performs half unordered compare not-equal operation.
__host__ __device__ bool operator<(const __half &lh, const __half &rh)
Performs half ordered less-than compare operation.
__host__ __device__ bool operator<=(const __half &lh, const __half &rh)
Performs half ordered less-or-equal compare operation.
__host__ __device__ bool operator==(const __half &lh, const __half &rh)
Performs half ordered compare equal operation.
__host__ __device__ bool operator>(const __half &lh, const __half &rh)
Performs half ordered greater-than compare operation.
__host__ __device__ bool operator>=(const __half &lh, const __half &rh)
Performs half ordered greater-or-equal compare operation.
4.3.1. Functions
__host__ __device__ bool __heq(const __half a, const __half b)
Performs half if-equal comparison.
Performs half if-equal comparison of inputs a and b. NaN inputs generate false results.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of if-equal comparison of a and b.
__host__ __device__ bool __hequ(const __half a, const __half b)
Performs half unordered if-equal comparison.
Performs half if-equal comparison of inputs a and b. NaN inputs generate true results.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of unordered if-equal comparison of a and b.
__host__ __device__ bool __hge(const __half a, const __half b)
Performs half greater-equal comparison.
Performs half greater-equal comparison of inputs a and b. NaN inputs generate false results.
Parameters
54
Chapter 4. Half Precision Intrinsics
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of greater-equal comparison of a and b.
__host__ __device__ bool __hgeu(const __half a, const __half b)
Performs half unordered greater-equal comparison.
Performs half greater-equal comparison of inputs a and b. NaN inputs generate true results.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of unordered greater-equal comparison of a and b.
__host__ __device__ bool __hgt(const __half a, const __half b)
Performs half greater-than comparison.
Performs half greater-than comparison of inputs a and b. NaN inputs generate false results.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of greater-than comparison of a and b.
__host__ __device__ bool __hgtu(const __half a, const __half b)
Performs half unordered greater-than comparison.
Performs half greater-than comparison of inputs a and b. NaN inputs generate true results.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of unordered greater-than comparison of a and b.
__host__ __device__ int __hisinf(const __half a)
Checks if the input half number is infinite.
Checks if the input half number a is infinite.
Parameters
a - [in] - half. Is only being read.
4.3. Half Comparison Functions
55
Returns
int
▶ -1 if a is equal to negative infinity,
▶ 1 if a is equal to positive infinity,
▶ 0 otherwise.
__host__ __device__ bool __hisnan(const __half a)
Determine whether half argument is a NaN.
Determine whether half value a is a NaN.
Parameters
a - [in] - half. Is only being read.
Returns
bool
▶ true if argument is NaN.
__host__ __device__ bool __hle(const __half a, const __half b)
Performs half less-equal comparison.
Performs half less-equal comparison of inputs a and b. NaN inputs generate false results.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of less-equal comparison of a and b.
__host__ __device__ bool __hleu(const __half a, const __half b)
Performs half unordered less-equal comparison.
Performs half less-equal comparison of inputs a and b. NaN inputs generate true results.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of unordered less-equal comparison of a and b.
__host__ __device__ bool __hlt(const __half a, const __half b)
Performs half less-than comparison.
Performs half less-than comparison of inputs a and b. NaN inputs generate false results.
Parameters
▶ a - [in] - half. Is only being read.
56
Chapter 4. Half Precision Intrinsics
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of less-than comparison of a and b.
__host__ __device__ bool __hltu(const __half a, const __half b)
Performs half unordered less-than comparison.
Performs half less-than comparison of inputs a and b. NaN inputs generate true results.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of unordered less-than comparison of a and b.
__host__ __device__ __half __hmax(const __half a, const __half b)
Calculates half maximum of two input values.
Calculates half max(a, b) defined as (a > b) ? a : b.
▶ If either of inputs is NaN, the other input is returned.
▶ If both inputs are NaNs, then canonical NaN is returned.
▶ If values of both inputs are 0.0, then +0.0 > -0.0
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
__host__ __device__ __half __hmax_nan(const __half a, const __half b)
Calculates half maximum of two input values, NaNs pass through.
Calculates half max(a, b) defined as (a > b) ? a : b.
▶ If either of inputs is NaN, then canonical NaN is returned.
▶ If values of both inputs are 0.0, then +0.0 > -0.0
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
4.3. Half Comparison Functions
57
__host__ __device__ __half __hmin(const __half a, const __half b)
Calculates half minimum of two input values.
Calculates half min(a, b) defined as (a < b) ? a : b.
▶ If either of inputs is NaN, the other input is returned.
▶ If both inputs are NaNs, then canonical NaN is returned.
▶ If values of both inputs are 0.0, then +0.0 > -0.0
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
__host__ __device__ __half __hmin_nan(const __half a, const __half b)
Calculates half minimum of two input values, NaNs pass through.
Calculates half min(a, b) defined as (a < b) ? a : b.
▶ If either of inputs is NaN, then canonical NaN is returned.
▶ If values of both inputs are 0.0, then +0.0 > -0.0
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half
__host__ __device__ bool __hne(const __half a, const __half b)
Performs half not-equal comparison.
Performs half not-equal comparison of inputs a and b. NaN inputs generate false results.
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of not-equal comparison of a and b.
__host__ __device__ bool __hneu(const __half a, const __half b)
Performs half unordered not-equal comparison.
Performs half not-equal comparison of inputs a and b. NaN inputs generate true results.
Parameters
58
Chapter 4. Half Precision Intrinsics
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
bool
▶ The boolean result of unordered not-equal comparison of a and b.
__host__ __device__ bool operator!=(const __half &lh, const __half &rh)
Performs half unordered compare not-equal operation.
See also:
__hneu(__half, __half)
__host__ __device__ bool operator<(const __half &lh, const __half &rh)
Performs half ordered less-than compare operation.
See also:
__hlt(__half, __half)
__host__ __device__ bool operator<=(const __half &lh, const __half &rh)
Performs half ordered less-or-equal compare operation.
See also:
__hle(__half, __half)
__host__ __device__ bool operator==(const __half &lh, const __half &rh)
Performs half ordered compare equal operation.
See also:
__heq(__half, __half)
__host__ __device__ bool operator>(const __half &lh, const __half &rh)
Performs half ordered greater-than compare operation.
See also:
__hgt(__half, __half)
__host__ __device__ bool operator>=(const __half &lh, const __half &rh)
Performs half ordered greater-or-equal compare operation.
See also:
__hge(__half, __half)
4.3. Half Comparison Functions
59
4.4. Half Math Functions
To use these functions, include the header file cuda_fp16.h in your program.
Functions
__device__ __half hceil(const __half h)
Calculate ceiling of the input argument.
__device__ __half hcos(const __half a)
Calculates half cosine in round-to-nearest-even mode.
__device__ __half hexp(const __half a)
Calculates half natural exponential function in round-to-nearest-even mode.
__device__ __half hexp10(const __half a)
Calculates half decimal exponential function in round-to-nearest-even mode.
__device__ __half hexp2(const __half a)
Calculates half binary exponential function in round-to-nearest-even mode.
__device__ __half hfloor(const __half h)
Calculate the largest integer less than or equal to h .
__device__ __half hlog(const __half a)
Calculates half natural logarithm in round-to-nearest-even mode.
__device__ __half hlog10(const __half a)
Calculates half decimal logarithm in round-to-nearest-even mode.
__device__ __half hlog2(const __half a)
Calculates half binary logarithm in round-to-nearest-even mode.
__device__ __half hrcp(const __half a)
Calculates half reciprocal in round-to-nearest-even mode.
__device__ __half hrint(const __half h)
Round input to nearest integer value in half-precision floating-point number.
__device__ __half hrsqrt(const __half a)
Calculates half reciprocal square root in round-to-nearest-even mode.
__device__ __half hsin(const __half a)
Calculates half sine in round-to-nearest-even mode.
__device__ __half hsqrt(const __half a)
Calculates half square root in round-to-nearest-even mode.
__device__ __half htanh(const __half a)
Calculates half hyperbolic tangent function in round-to-nearest-even mode.
__device__ __half htanh_approx(const __half a)
Calculates approximate half hyperbolic tangent function.
__device__ __half htrunc(const __half h)
Truncate input argument to the integral part.
60
Chapter 4. Half Precision Intrinsics
4.4.1. Functions
__device__ __half hceil(const __half h)
Calculate ceiling of the input argument.
Compute the smallest integer value not less than h.
Parameters
h - [in] - half. Is only being read.
Returns
half
▶ The smallest integer value not less than h.
▶ hceil( ±0 ) returns ±0.
▶ hceil( ±∞ ) returns ±∞.
▶ hceil(NaN) returns NaN.
__device__ __half hcos(const __half a)
Calculates half cosine in round-to-nearest-even mode.
Calculates half cosine of input a in round-to-nearest-even mode.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The cosine of a.
▶ hcos (±0) returns 1.
▶ hcos (±∞) returns NaN.
▶ hcos(NaN) returns NaN.
__device__ __half hexp(const __half a)
Calculates half natural exponential function in round-to-nearest-even mode.
Calculates half natural exponential function of input: ea in round-to-nearest-even mode.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The natural exponential function on a.
▶ hexp (±0) returns 1.
▶ hexp (-∞) returns +0.
▶ hexp (+∞) returns +∞.
▶ hexp(NaN) returns NaN.
__device__ __half hexp10(const __half a)
Calculates half decimal exponential function in round-to-nearest-even mode.
Calculates half decimal exponential function of input: 10a in round-to-nearest-even mode.
4.4. Half Math Functions
61
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The decimal exponential function on a.
▶ hexp10 (±0) returns 1.
▶ hexp10 (-∞) returns +0.
▶ hexp10 (+∞) returns +∞.
▶ hexp10(NaN) returns NaN.
__device__ __half hexp2(const __half a)
Calculates half binary exponential function in round-to-nearest-even mode.
Calculates half binary exponential function of input: 2a in round-to-nearest-even mode.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The binary exponential function on a.
▶ hexp2 (±0) returns 1.
▶ hexp2 (-∞) returns +0.
▶ hexp2 (+∞) returns +∞.
▶ hexp2(NaN) returns NaN.
__device__ __half hfloor(const __half h)
Calculate the largest integer less than or equal to h.
Calculate the largest integer value which is less than or equal to h.
Parameters
h - [in] - half. Is only being read.
Returns
half
▶ The largest integer value which is less than or equal to h.
▶ hfloor( ±0 ) returns ±0.
▶ hfloor( ±∞ ) returns ±∞.
▶ hfloor(NaN) returns NaN.
__device__ __half hlog(const __half a)
Calculates half natural logarithm in round-to-nearest-even mode.
Calculates half natural logarithm of input: ln(a) in round-to-nearest-even mode.
Parameters
a - [in] - half. Is only being read.
Returns
half
62
Chapter 4. Half Precision Intrinsics
▶ The natural logarithm of a.
▶ hlog (±0) returns -∞.
▶ hlog(1) returns +0.
▶ hlog(x), x < 0 returns NaN.
▶ hlog (+∞) returns +∞.
▶ hlog(NaN) returns NaN.
__device__ __half hlog10(const __half a)
Calculates half decimal logarithm in round-to-nearest-even mode.
Calculates half decimal logarithm of input: log10(a) in round-to-nearest-even mode.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The decimal logarithm of a.
▶ hlog10 (±0) returns -∞.
▶ hlog10(1) returns +0.
▶ hlog10(x), x < 0 returns NaN.
▶ hlog10 (+∞) returns +∞.
▶ hlog10(NaN) returns NaN.
__device__ __half hlog2(const __half a)
Calculates half binary logarithm in round-to-nearest-even mode.
Calculates half binary logarithm of input: log2(a) in round-to-nearest-even mode.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The binary logarithm of a.
▶ hlog2 (±0) returns -∞.
▶ hlog2(1) returns +0.
▶ hlog2(x), x < 0 returns NaN.
▶ hlog2 (+∞) returns +∞.
▶ hlog2(NaN) returns NaN.
__device__ __half hrcp(const __half a)
Calculates half reciprocal in round-to-nearest-even mode.
Calculates half reciprocal of input:1a in round-to-nearest-even mode.
Parameters
a - [in] - half. Is only being read.
Returns
half
4.4. Half Math Functions
63
▶ The reciprocal of a.
▶ hrcp (±0) returns ±∞.
▶ hrcp (±∞) returns ±0.
▶ hrcp(NaN) returns NaN.
__device__ __half hrint(const __half h)
Round input to nearest integer value in half-precision floating-point number.
Round h to the nearest integer value in half-precision floating-point format, with halfway cases
rounded to the nearest even integer value.
Parameters
h - [in] - half. Is only being read.
Returns
half
▶ The nearest integer to h.
▶ hrint( ±0 ) returns ±0.
▶ hrint( ±∞ ) returns ±∞.
▶ hrint(NaN) returns NaN.
__device__ __half hrsqrt(const __half a)
Calculates half reciprocal square root in round-to-nearest-even mode.
Calculates half reciprocal square root of input:
1
√
in round-to-nearest-even mode.
a
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The reciprocal square root of a.
▶ hrsqrt (±0) returns ±∞.
▶ hrsqrt (+∞) returns +0.
▶ hrsqrt (x), x < 0.0 returns NaN.
▶ hrsqrt(NaN) returns NaN.
__device__ __half hsin(const __half a)
Calculates half sine in round-to-nearest-even mode.
Calculates half sine of input a in round-to-nearest-even mode.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The sine of a.
▶ hsin (±0) returns (±0).
▶ hsin (±∞) returns NaN.
▶ hsin(NaN) returns NaN.
64
Chapter 4. Half Precision Intrinsics
__device__ __half hsqrt(const __half a)
Calculates half square root in round-to-nearest-even mode.
Calculates half square root of input:
√a in round-to-nearest-even mode.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The square root of a.
▶ hsqrt (+∞) returns +∞.
▶ hsqrt (±0) returns ±0.
▶ hsqrt (x), x < 0.0 returns NaN.
▶ hsqrt(NaN) returns NaN.
__device__ __half htanh(const __half a)
Calculates half hyperbolic tangent function in round-to-nearest-even mode.
Calculates half hyperbolic tangent function: tanh(a) in round-to-nearest-even mode.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The hyperbolic tangent function of a.
▶ htanh (±0) returns (±0).
▶ htanh (±∞) returns (±1).
▶ htanh(NaN) returns NaN.
__device__ __half htanh_approx(const __half a)
Calculates approximate half hyperbolic tangent function.
Calculates approximate half hyperbolic tangent function: tanh(a). This operation uses HW ac-
celeration on devices of compute capability 7.5 and higher.
Parameters
a - [in] - half. Is only being read.
Returns
half
▶ The approximate hyperbolic tangent function of a.
▶ htanh_approx (±0) returns (±0).
▶ htanh_approx (±∞) returns (±1).
▶ htanh_approx(NaN) returns NaN.
__device__ __half htrunc(const __half h)
Truncate input argument to the integral part.
Round h to the largest integer value that does not exceed h in magnitude.
Parameters
h - [in] - half. Is only being read.
4.4. Half Math Functions
65
Returns
half
▶ The truncated value.
▶ htrunc( ±0 ) returns ±0.
▶ htrunc( ±∞ ) returns ±∞.
▶ htrunc(NaN) returns NaN.
4.5. Half Precision Conversion and Data
Movement
To use these functions, include the header file cuda_fp16.h in your program.
Functions
__host__ __device__ __half __double2half(const double a)
Converts double number to half precision in round-to-nearest-even mode and returns half with
converted value.
__host__ __device__ __half2 __float22half2_rn(const float2 a)
Converts both components of float2 number to half precision in round-to-nearest-even mode
and returns half2 with converted values.
__host__ __device__ __half __float2half(const float a)
Converts float number to half precision in round-to-nearest-even mode and returns half with
converted value.
__host__ __device__ __half2 __float2half2_rn(const float a)
Converts input to half precision in round-to-nearest-even mode and populates both halves of
half2 with converted value.
__host__ __device__ __half __float2half_rd(const float a)
Converts float number to half precision in round-down mode and returns half with converted
value.
__host__ __device__ __half __float2half_rn(const float a)
Converts float number to half precision in round-to-nearest-even mode and returns half with
converted value.
__host__ __device__ __half __float2half_ru(const float a)
Converts float number to half precision in round-up mode and returns half with converted value.
__host__ __device__ __half __float2half_rz(const float a)
Converts float number to half precision in round-towards-zero mode and returns half with con-
verted value.
__host__ __device__ __half2 __floats2half2_rn(const float a, const float b)
Converts both input floats to half precision in round-to-nearest-even mode and returns half2
with converted values.
__host__ __device__ float2 __half22float2(const __half2 a)
Converts both halves of half2 to float2 and returns the result.
66
Chapter 4. Half Precision Intrinsics
__host__ __device__ __half2::__half2(const __half2_raw &h2r)
Constructor from __half2_raw .
__host__ __device__ constexpr __half2::__half2(const __half &a, const __half &b)
Constructor from two __half variables.
__host__ __device__ __half2::__half2(const __half2 &&src)
Move constructor, available for C++11 and later dialects.
__host__ __device__ __half2::__half2(const __half2 &src)
Copy constructor.
__half2::__half2()=default
Constructor by default.
__host__ __device__ __half2::operator __half2_raw() const
Conversion operator to __half2_raw .
__host__ __device__ __half2 & __half2::operator=(const __half2_raw &h2r)
Assignment operator from __half2_raw .
__host__ __device__ __half2 & __half2::operator=(const __half2 &&src)
Move assignment operator, available for C++11 and later dialects.
__host__ __device__ __half2 & __half2::operator=(const __half2 &src)
Copy assignment operator.
__host__ __device__ signed char __half2char_rz(const __half h)
Convert a half to a signed char in round-towards-zero mode.
__host__ __device__ float __half2float(const __half a)
Converts half number to float.
__host__ __device__ __half2 __half2half2(const __half a)
Returns half2 with both halves equal to the input value.
__device__ int __half2int_rd(const __half h)
Convert a half to a signed integer in round-down mode.
__device__ int __half2int_rn(const __half h)
Convert a half to a signed integer in round-to-nearest-even mode.
__device__ int __half2int_ru(const __half h)
Convert a half to a signed integer in round-up mode.
__host__ __device__ int __half2int_rz(const __half h)
Convert a half to a signed integer in round-towards-zero mode.
__device__ long long int __half2ll_rd(const __half h)
Convert a half to a signed 64-bit integer in round-down mode.
__device__ long long int __half2ll_rn(const __half h)
Convert a half to a signed 64-bit integer in round-to-nearest-even mode.
__device__ long long int __half2ll_ru(const __half h)
Convert a half to a signed 64-bit integer in round-up mode.
__host__ __device__ long long int __half2ll_rz(const __half h)
Convert a half to a signed 64-bit integer in round-towards-zero mode.
__device__ short int __half2short_rd(const __half h)
Convert a half to a signed short integer in round-down mode.
4.5. Half Precision Conversion and Data Movement
67
__device__ short int __half2short_rn(const __half h)
Convert a half to a signed short integer in round-to-nearest-even mode.
__device__ short int __half2short_ru(const __half h)
Convert a half to a signed short integer in round-up mode.
__host__ __device__ short int __half2short_rz(const __half h)
Convert a half to a signed short integer in round-towards-zero mode.
__host__ __device__ unsigned char __half2uchar_rz(const __half h)
Convert a half to an unsigned char in round-towards-zero mode.
__device__ unsigned int __half2uint_rd(const __half h)
Convert a half to an unsigned integer in round-down mode.
__device__ unsigned int __half2uint_rn(const __half h)
Convert a half to an unsigned integer in round-to-nearest-even mode.
__device__ unsigned int __half2uint_ru(const __half h)
Convert a half to an unsigned integer in round-up mode.
__host__ __device__ unsigned int __half2uint_rz(const __half h)
Convert a half to an unsigned integer in round-towards-zero mode.
__device__ unsigned long long int __half2ull_rd(const __half h)
Convert a half to an unsigned 64-bit integer in round-down mode.
__device__ unsigned long long int __half2ull_rn(const __half h)
Convert a half to an unsigned 64-bit integer in round-to-nearest-even mode.
__device__ unsigned long long int __half2ull_ru(const __half h)
Convert a half to an unsigned 64-bit integer in round-up mode.
__host__ __device__ unsigned long long int __half2ull_rz(const __half h)
Convert a half to an unsigned 64-bit integer in round-towards-zero mode.
__device__ unsigned short int __half2ushort_rd(const __half h)
Convert a half to an unsigned short integer in round-down mode.
__device__ unsigned short int __half2ushort_rn(const __half h)
Convert a half to an unsigned short integer in round-to-nearest-even mode.
__device__ unsigned short int __half2ushort_ru(const __half h)
Convert a half to an unsigned short integer in round-up mode.
__host__ __device__ unsigned short int __half2ushort_rz(const __half h)
Convert a half to an unsigned short integer in round-towards-zero mode.
__host__ __device__ constexpr __half::__half(const __half_raw &hr)
Constructor from __half_raw .
__host__ __device__ __half::__half(const unsigned short val)
Construct __half from unsigned short integer input using default round-to-nearest-even
rounding mode.
__host__ __device__ __half::__half(const unsigned int val)
Construct __half from unsigned int input using default round-to-nearest-even rounding
mode.
__host__ __device__ __half::__half(const short val)
Construct __half from short integer input using default round-to-nearest-even rounding
mode.
68
Chapter 4. Half Precision Intrinsics
__host__ __device__ __half::__half(const double f)
Construct __half from double input using default round-to-nearest-even rounding mode.
__host__ __device__ __half::__half(const unsigned long val)
Construct __half from unsigned long input using default round-to-nearest-even rounding
mode.
__host__ __device__ __half::__half(const float f)
Construct __half from float input using default round-to-nearest-even rounding mode.
__host__ __device__ __half::__half(const int val)
Construct __half from int input using default round-to-nearest-even rounding mode.
__host__ __device__ __half::__half(const long val)
Construct __half from long input using default round-to-nearest-even rounding mode.
__host__ __device__ __half::__half(const long long val)
Construct __half from long long input using default round-to-nearest-even rounding mode.
__half::__half()=default
Constructor by default.
__host__ __device__ __half::__half(const __nv_bfloat16 f)
Construct __half from __nv_bfloat16 input using default round-to-nearest-even rounding
mode.
__host__ __device__ __half::__half(const unsigned long long val)
Construct __half from unsigned long long input using default round-to-nearest-even round-
ing mode.
__host__ __device__ __half::operator __half_raw() const volatile
Type cast to __half_raw operator with volatile input.
__host__ __device__ __half::operator __half_raw() const
Type cast to __half_raw operator.
__host__ __device__ constexpr __half::operator bool() const
Conversion operator to bool data type.
__host__ __device__ __half::operator char() const
Conversion operator to an implementation defined char data type.
__host__ __device__ __half::operator float() const
Type cast to float operator.
__host__ __device__ __half::operator int() const
Conversion operator to int data type.
__host__ __device__ __half::operator long() const
Conversion operator to long data type.
__host__ __device__ __half::operator long long() const
Conversion operator to long long data type.
__host__ __device__ __half::operator short() const
Conversion operator to short data type.
__host__ __device__ __half::operator signed char() const
Conversion operator to signed char data type.
__host__ __device__ __half::operator unsigned char() const
Conversion operator to unsigned char data type.
4.5. Half Precision Conversion and Data Movement
69
__host__ __device__ __half::operator unsigned int() const
Conversion operator to unsigned int data type.
__host__ __device__ __half::operator unsigned long() const
Conversion operator to unsigned long data type.
__host__ __device__ __half::operator unsigned long long() const
Conversion operator to unsigned long long data type.
__host__ __device__ __half::operator unsigned short() const
Conversion operator to unsigned short data type.
__host__ __device__ __half & __half::operator=(const float f)
Type cast to __half assignment operator from float input using default round-to-nearest-even
rounding mode.
__host__ __device__ volatile __half & __half::operator=(const volatile __half_raw &hr) volatile
Assignment operator from volatile __half_raw to volatile __half .
__host__ __device__ __half & __half::operator=(const long long val)
Type cast from long long assignment operator, using default round-to-nearest-even rounding
mode.
__host__ __device__ volatile __half & __half::operator=(const __half_raw &hr) volatile
Assignment operator from __half_raw to volatile __half .
__host__ __device__ __half & __half::operator=(const unsigned int val)
Type cast from unsigned int assignment operator, using default round-to-nearest-even round-
ing mode.
__host__ __device__ __half & __half::operator=(const unsigned short val)
Type cast from unsigned short assignment operator, using default round-to-nearest-even
rounding mode.
__host__ __device__ __half & __half::operator=(const short val)
Type cast from short assignment operator, using default round-to-nearest-even rounding mode.
__host__ __device__ __half & __half::operator=(const double f)
Type cast to __half assignment operator from double input using default round-to-nearest-
even rounding mode.
__host__ __device__ __half & __half::operator=(const __half_raw &hr)
Assignment operator from __half_raw .
__host__ __device__ __half & __half::operator=(const unsigned long long val)
Type cast from unsigned long long assignment operator, using default round-to-nearest-even
rounding mode.
__host__ __device__ __half & __half::operator=(const int val)
Type cast from int assignment operator, using default round-to-nearest-even rounding mode.
__host__ __device__ short int __half_as_short(const __half h)
Reinterprets bits in a half as a signed short integer.
__host__ __device__ unsigned short int __half_as_ushort(const __half h)
Reinterprets bits in a half as an unsigned short integer.
__host__ __device__ __half2 __halves2half2(const __half a, const __half b)
Combines two half numbers into one half2 number.
__host__ __device__ float __high2float(const __half2 a)
Converts high 16 bits of half2 to float and returns the result.
70
Chapter 4. Half Precision Intrinsics
__host__ __device__ __half __high2half(const __half2 a)
Returns high 16 bits of half2 input.
__host__ __device__ __half2 __high2half2(const __half2 a)
Extracts high 16 bits from half2 input.
__host__ __device__ __half2 __highs2half2(const __half2 a, const __half2 b)
Extracts high 16 bits from each of the two half2 inputs and combines into one half2 number.
__host__ __device__ __half __int2half_rd(const int i)
Convert a signed integer to a half in round-down mode.
__host__ __device__ __half __int2half_rn(const int i)
Convert a signed integer to a half in round-to-nearest-even mode.
__host__ __device__ __half __int2half_ru(const int i)
Convert a signed integer to a half in round-up mode.
__host__ __device__ __half __int2half_rz(const int i)
Convert a signed integer to a half in round-towards-zero mode.
__device__ __half2 __ldca(const __half2 *const ptr)
Generates a ld.global.ca load instruction.
__device__ __half __ldca(const __half *const ptr)
Generates a ld.global.ca load instruction.
__device__ __half __ldcg(const __half *const ptr)
Generates a ld.global.cg load instruction.
__device__ __half2 __ldcg(const __half2 *const ptr)
Generates a ld.global.cg load instruction.
__device__ __half __ldcs(const __half *const ptr)
Generates a ld.global.cs load instruction.
__device__ __half2 __ldcs(const __half2 *const ptr)
Generates a ld.global.cs load instruction.
__device__ __half2 __ldcv(const __half2 *const ptr)
Generates a ld.global.cv load instruction.
__device__ __half __ldcv(const __half *const ptr)
Generates a ld.global.cv load instruction.
__device__ __half2 __ldg(const __half2 *const ptr)
Generates a ld.global.nc load instruction.
__device__ __half __ldg(const __half *const ptr)
Generates a ld.global.nc load instruction.
__device__ __half __ldlu(const __half *const ptr)
Generates a ld.global.lu load instruction.
__device__ __half2 __ldlu(const __half2 *const ptr)
Generates a ld.global.lu load instruction.
__host__ __device__ __half __ll2half_rd(const long long int i)
Convert a signed 64-bit integer to a half in round-down mode.
__host__ __device__ __half __ll2half_rn(const long long int i)
Convert a signed 64-bit integer to a half in round-to-nearest-even mode.
4.5. Half Precision Conversion and Data Movement
71
__host__ __device__ __half __ll2half_ru(const long long int i)
Convert a signed 64-bit integer to a half in round-up mode.
__host__ __device__ __half __ll2half_rz(const long long int i)
Convert a signed 64-bit integer to a half in round-towards-zero mode.
__host__ __device__ float __low2float(const __half2 a)
Converts low 16 bits of half2 to float and returns the result.
__host__ __device__ __half __low2half(const __half2 a)
Returns low 16 bits of half2 input.
__host__ __device__ __half2 __low2half2(const __half2 a)
Extracts low 16 bits from half2 input.
__host__ __device__ __half2 __lowhigh2highlow(const __half2 a)
Swaps both halves of the half2 input.
__host__ __device__ __half2 __lows2half2(const __half2 a, const __half2 b)
Extracts low 16 bits from each of the two half2 inputs and combines into one half2 number.
__device__ __half __shfl_down_sync(const unsigned int mask, const __half var, const unsigned int
delta, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __half2 __shfl_down_sync(const unsigned int mask, const __half2 var, const unsigned int
delta, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __half2 __shfl_sync(const unsigned int mask, const __half2 var, const int srcLane, const
int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __half __shfl_sync(const unsigned int mask, const __half var, const int srcLane, const int
width=warpSize)
Exchange a variable between threads within a warp.
__device__ __half2 __shfl_up_sync(const unsigned int mask, const __half2 var, const unsigned int
delta, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __half __shfl_up_sync(const unsigned int mask, const __half var, const unsigned int delta,
const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __half2 __shfl_xor_sync(const unsigned int mask, const __half2 var, const int laneMask,
const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __half __shfl_xor_sync(const unsigned int mask, const __half var, const int laneMask,
const int width=warpSize)
Exchange a variable between threads within a warp.
__host__ __device__ __half __short2half_rd(const short int i)
Convert a signed short integer to a half in round-down mode.
__host__ __device__ __half __short2half_rn(const short int i)
Convert a signed short integer to a half in round-to-nearest-even mode.
__host__ __device__ __half __short2half_ru(const short int i)
Convert a signed short integer to a half in round-up mode.
72
Chapter 4. Half Precision Intrinsics
__host__ __device__ __half __short2half_rz(const short int i)
Convert a signed short integer to a half in round-towards-zero mode.
__host__ __device__ __half __short_as_half(const short int i)
Reinterprets bits in a signed short integer as a half .
__device__ void __stcg(__half2 *const ptr, const __half2 value)
Generates a st.global.cg store instruction.
__device__ void __stcg(__half *const ptr, const __half value)
Generates a st.global.cg store instruction.
__device__ void __stcs(__half2 *const ptr, const __half2 value)
Generates a st.global.cs store instruction.
__device__ void __stcs(__half *const ptr, const __half value)
Generates a st.global.cs store instruction.
__device__ void __stwb(__half2 *const ptr, const __half2 value)
Generates a st.global.wb store instruction.
__device__ void __stwb(__half *const ptr, const __half value)
Generates a st.global.wb store instruction.
__device__ void __stwt(__half *const ptr, const __half value)
Generates a st.global.wt store instruction.
__device__ void __stwt(__half2 *const ptr, const __half2 value)
Generates a st.global.wt store instruction.
__host__ __device__ __half __uint2half_rd(const unsigned int i)
Convert an unsigned integer to a half in round-down mode.
__host__ __device__ __half __uint2half_rn(const unsigned int i)
Convert an unsigned integer to a half in round-to-nearest-even mode.
__host__ __device__ __half __uint2half_ru(const unsigned int i)
Convert an unsigned integer to a half in round-up mode.
__host__ __device__ __half __uint2half_rz(const unsigned int i)
Convert an unsigned integer to a half in round-towards-zero mode.
__host__ __device__ __half __ull2half_rd(const unsigned long long int i)
Convert an unsigned 64-bit integer to a half in round-down mode.
__host__ __device__ __half __ull2half_rn(const unsigned long long int i)
Convert an unsigned 64-bit integer to a half in round-to-nearest-even mode.
__host__ __device__ __half __ull2half_ru(const unsigned long long int i)
Convert an unsigned 64-bit integer to a half in round-up mode.
__host__ __device__ __half __ull2half_rz(const unsigned long long int i)
Convert an unsigned 64-bit integer to a half in round-towards-zero mode.
__host__ __device__ __half __ushort2half_rd(const unsigned short int i)
Convert an unsigned short integer to a half in round-down mode.
__host__ __device__ __half __ushort2half_rn(const unsigned short int i)
Convert an unsigned short integer to a half in round-to-nearest-even mode.
__host__ __device__ __half __ushort2half_ru(const unsigned short int i)
Convert an unsigned short integer to a half in round-up mode.
4.5. Half Precision Conversion and Data Movement
73
__host__ __device__ __half __ushort2half_rz(const unsigned short int i)
Convert an unsigned short integer to a half in round-towards-zero mode.
__host__ __device__ __half __ushort_as_half(const unsigned short int i)
Reinterprets bits in an unsigned short integer as a half .
__host__ __device__ __half2 make_half2(const __half x, const __half y)
Vector function, combines two __half numbers into one __half2 number.
4.5.1. Functions
__host__ __device__ __half __double2half(const double a)
Converts double number to half precision in round-to-nearest-even mode and returns half with
converted value.
Converts double number a to half precision in round-to-nearest-even mode.
Parameters
a - [in] - double. Is only being read.
Returns
half
▶ a converted to half precision using round-to-nearest-even mode.
▶ __double2half (±0) returns ±0.
▶ __double2half (±∞) returns ±∞.
▶ __double2half(NaN) returns NaN.
__host__ __device__ __half2 __float22half2_rn(const float2 a)
Converts both components of float2 number to half precision in round-to-nearest-even mode
and returns half2 with converted values.
Converts both components of float2 to half precision in round-to-nearest-even mode and com-
bines the results into one half2 number. Low 16 bits of the return value correspond to a.x and
high 16 bits of the return value correspond to a.y.
See also:
__float2half_rn(float) for further details.
Parameters
a - [in] - float2. Is only being read.
Returns
half2
▶ The half2 which has corresponding halves equal to the converted float2
components.
__host__ __device__ __half __float2half(const float a)
Converts float number to half precision in round-to-nearest-even mode and returns half with
converted value.
Converts float number a to half precision in round-to-nearest-even mode.
74
Chapter 4. Half Precision Intrinsics
////////////////////////////////////////// |
||
|
|
|