|
|
See also:
__float2half_rn(float) for further details.
Parameters
a - [in] - float. Is only being read.
Returns
half
▶ a converted to half precision using round-to-nearest-even mode.
__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.
Converts input a to half precision in round-to-nearest-even mode and populates both halves of
half2 with converted value.
See also:
__float2half_rn(float) for further details.
Parameters
a - [in] - float. Is only being read.
Returns
half2
▶ The half2 value with both halves equal to the converted half precision number.
__host__ __device__ __half __float2half_rd(const float a)
Converts float number to half precision in round-down mode and returns half with converted
value.
Converts float number a to half precision in round-down mode.
Parameters
a - [in] - float. Is only being read.
Returns
half
▶ a converted to half precision using round-down mode.
▶ __float2half_rd (±0) returns ±0.
▶ __float2half_rd (±∞) returns ±∞.
▶ __float2half_rd(NaN) returns NaN.
__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.
Converts float number a to half precision in round-to-nearest-even mode.
Parameters
a - [in] - float. Is only being read.
Returns
half
4.5. Half Precision Conversion and Data Movement
75
▶ a converted to half precision using round-to-nearest-even mode.
▶ __float2half_rn (±0) returns ±0.
▶ __float2half_rn (±∞) returns ±∞.
▶ __float2half_rn(NaN) returns NaN.
__host__ __device__ __half __float2half_ru(const float a)
Converts float number to half precision in round-up mode and returns half with converted value.
Converts float number a to half precision in round-up mode.
Parameters
a - [in] - float. Is only being read.
Returns
half
▶ a converted to half precision using round-up mode.
▶ __float2half_ru (±0) returns ±0.
▶ __float2half_ru (±∞) returns ±∞.
▶ __float2half_ru(NaN) returns NaN.
__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.
Converts float number a to half precision in round-towards-zero mode.
Parameters
a - [in] - float. Is only being read.
Returns
half
▶ a converted to half precision using round-towards-zero mode.
▶ __float2half_rz (±0) returns ±0.
▶ __float2half_rz (±∞) returns ±∞.
▶ __float2half_rz(NaN) returns NaN.
__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.
Converts both input floats to half precision in round-to-nearest-even mode and combines the
results into one half2 number. Low 16 bits of the return value correspond to the input a, high
16 bits correspond to the input b.
See also:
__float2half_rn(float) for further details.
Parameters
▶ a - [in] - float. Is only being read.
▶ b - [in] - float. Is only being read.
76
Chapter 4. Half Precision Intrinsics
Returns
half2
▶ The half2 value with corresponding halves equal to the converted input floats.
__host__ __device__ float2 __half22float2(const __half2 a)
Converts both halves of half2 to float2 and returns the result.
Converts both halves of half2 input a to float2 and returns the result.
See also:
__half2float(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
float2
▶ a converted to float2.
__host__ __device__ signed char __half2char_rz(const __half h)
Convert a half to a signed char in round-towards-zero mode.
Convert the half-precision floating-point value h to a signed char integer in round-towards-zero
mode. NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
signed char
▶ h converted to a signed char using round-towards-zero mode.
▶ __half2char_rz (±0) returns 0.
▶ __half2char_rz (x), x > 127 returns SCHAR_MAX = 0x7F.
▶ __half2char_rz (x), x < -128 returns SCHAR_MIN = 0x80.
▶ __half2char_rz(NaN) returns 0.
__host__ __device__ float __half2float(const __half a)
Converts half number to float.
Converts half number a to float.
Parameters
a - [in] - float. Is only being read.
Returns
float
▶ a converted to float.
▶ __half2float (±0) returns ±0.
▶ __half2float (±∞) returns ±∞.
▶ __half2float(NaN) returns NaN.
4.5. Half Precision Conversion and Data Movement
77
__host__ __device__ __half2 __half2half2(const __half a)
Returns half2 with both halves equal to the input value.
Returns half2 number with both halves equal to the input a half number.
Parameters
a - [in] - half. Is only being read.
Returns
half2
▶ The vector which has both its halves equal to the input a.
__device__ int __half2int_rd(const __half h)
Convert a half to a signed integer in round-down mode.
Convert the half-precision floating-point value h to a signed integer in round-down mode. NaN
inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
int
▶ h converted to a signed integer using round-down mode.
▶ __half2int_rd (±0) returns 0.
▶ __half2int_rd (+∞) returns INT_MAX = 0x7FFFFFFF.
▶ __half2int_rd (-∞) returns INT_MIN = 0x80000000.
▶ __half2int_rd(NaN) returns 0.
__device__ int __half2int_rn(const __half h)
Convert a half to a signed integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to a signed integer in round-to-nearest-even
mode. NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
int
▶ h converted to a signed integer using round-to-nearest-even mode.
▶ __half2int_rn (±0) returns 0.
▶ __half2int_rn (+∞) returns INT_MAX = 0x7FFFFFFF.
▶ __half2int_rn (-∞) returns INT_MIN = 0x80000000.
▶ __half2int_rn(NaN) returns 0.
__device__ int __half2int_ru(const __half h)
Convert a half to a signed integer in round-up mode.
Convert the half-precision floating-point value h to a signed integer in round-up mode. NaN
inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
78
Chapter 4. Half Precision Intrinsics
Returns
int
▶ h converted to a signed integer using round-up mode.
▶ __half2int_ru (±0) returns 0.
▶ __half2int_ru (+∞) returns INT_MAX = 0x7FFFFFFF.
▶ __half2int_ru (-∞) returns INT_MIN = 0x80000000.
▶ __half2int_ru(NaN) returns 0.
__host__ __device__ int __half2int_rz(const __half h)
Convert a half to a signed integer in round-towards-zero mode.
Convert the half-precision floating-point value h to a signed integer in round-towards-zero mode.
NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
int
▶ h converted to a signed integer using round-towards-zero mode.
▶ __half2int_rz (±0) returns 0.
▶ __half2int_rz (+∞) returns INT_MAX = 0x7FFFFFFF.
▶ __half2int_rz (-∞) returns INT_MIN = 0x80000000.
▶ __half2int_rz(NaN) returns 0.
__device__ long long int __half2ll_rd(const __half h)
Convert a half to a signed 64-bit integer in round-down mode.
Convert the half-precision floating-point value h to a signed 64-bit integer in round-down mode.
NaN inputs return a long long int with hex value of 0x8000000000000000.
Parameters
h - [in] - half. Is only being read.
Returns
long long int
▶ h converted to a signed 64-bit integer using round-down mode.
▶ __half2ll_rd (±0) returns 0.
▶ __half2ll_rd (+∞) returns LLONG_MAX = 0x7FFFFFFFFFFFFFFF.
▶ __half2ll_rd (-∞) returns LLONG_MIN = 0x8000000000000000.
▶ __half2ll_rd(NaN) returns 0x8000000000000000.
__device__ long long int __half2ll_rn(const __half h)
Convert a half to a signed 64-bit integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to a signed 64-bit integer in round-to-nearest-
even mode. NaN inputs return a long long int with hex value of 0x8000000000000000.
Parameters
h - [in] - half. Is only being read.
4.5. Half Precision Conversion and Data Movement
79
Returns
long long int
▶ h converted to a signed 64-bit integer using round-to-nearest-even mode.
▶ __half2ll_rn (±0) returns 0.
▶ __half2ll_rn (+∞) returns LLONG_MAX = 0x7FFFFFFFFFFFFFFF.
▶ __half2ll_rn (-∞) returns LLONG_MIN = 0x8000000000000000.
▶ __half2ll_rn(NaN) returns 0x8000000000000000.
__device__ long long int __half2ll_ru(const __half h)
Convert a half to a signed 64-bit integer in round-up mode.
Convert the half-precision floating-point value h to a signed 64-bit integer in round-up mode.
NaN inputs return a long long int with hex value of 0x8000000000000000.
Parameters
h - [in] - half. Is only being read.
Returns
long long int
▶ h converted to a signed 64-bit integer using round-up mode.
▶ __half2ll_ru (±0) returns 0.
▶ __half2ll_ru (+∞) returns LLONG_MAX = 0x7FFFFFFFFFFFFFFF.
▶ __half2ll_ru (-∞) returns LLONG_MIN = 0x8000000000000000.
▶ __half2ll_ru(NaN) returns 0x8000000000000000.
__host__ __device__ long long int __half2ll_rz(const __half h)
Convert a half to a signed 64-bit integer in round-towards-zero mode.
Convert the half-precision floating-point value h to a signed 64-bit integer in round-towards-zero
mode. NaN inputs return a long long int with hex value of 0x8000000000000000.
Parameters
h - [in] - half. Is only being read.
Returns
long long int
▶ h converted to a signed 64-bit integer using round-towards-zero mode.
▶ __half2ll_rz (±0) returns 0.
▶ __half2ll_rz (+∞) returns LLONG_MAX = 0x7FFFFFFFFFFFFFFF.
▶ __half2ll_rz (-∞) returns LLONG_MIN = 0x8000000000000000.
▶ __half2ll_rz(NaN) returns 0x8000000000000000.
__device__ short int __half2short_rd(const __half h)
Convert a half to a signed short integer in round-down mode.
Convert the half-precision floating-point value h to a signed short integer in round-down mode.
NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
80
Chapter 4. Half Precision Intrinsics
Returns
short int
▶ h converted to a signed short integer using round-down mode.
▶ __half2short_rd (±0) returns 0.
▶ __half2short_rd (x), x > 32767 returns SHRT_MAX = 0x7FFF.
▶ __half2short_rd (x), x < -32768 returns SHRT_MIN = 0x8000.
▶ __half2short_rd(NaN) returns 0.
__device__ short int __half2short_rn(const __half h)
Convert a half to a signed short integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to a signed short integer in round-to-nearest-
even mode. NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
short int
▶ h converted to a signed short integer using round-to-nearest-even mode.
▶ __half2short_rn (±0) returns 0.
▶ __half2short_rn (x), x > 32767 returns SHRT_MAX = 0x7FFF.
▶ __half2short_rn (x), x < -32768 returns SHRT_MIN = 0x8000.
▶ __half2short_rn(NaN) returns 0.
__device__ short int __half2short_ru(const __half h)
Convert a half to a signed short integer in round-up mode.
Convert the half-precision floating-point value h to a signed short integer in round-up mode. NaN
inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
short int
▶ h converted to a signed short integer using round-up mode.
▶ __half2short_ru (±0) returns 0.
▶ __half2short_ru (x), x > 32767 returns SHRT_MAX = 0x7FFF.
▶ __half2short_ru (x), x < -32768 returns SHRT_MIN = 0x8000.
▶ __half2short_ru(NaN) returns 0.
__host__ __device__ short int __half2short_rz(const __half h)
Convert a half to a signed short integer in round-towards-zero mode.
Convert the half-precision floating-point value h to a signed short integer in round-towards-zero
mode. NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
4.5. Half Precision Conversion and Data Movement
81
Returns
short int
▶ h converted to a signed short integer using round-towards-zero mode.
▶ __half2short_rz (±0) returns 0.
▶ __half2short_rz (x), x > 32767 returns SHRT_MAX = 0x7FFF.
▶ __half2short_rz (x), x < -32768 returns SHRT_MIN = 0x8000.
▶ __half2short_rz(NaN) returns 0.
__host__ __device__ unsigned char __half2uchar_rz(const __half h)
Convert a half to an unsigned char in round-towards-zero mode.
Convert the half-precision floating-point value h to an unsigned char in round-towards-zero
mode. NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
unsigned char
▶ h converted to an unsigned char using round-towards-zero mode.
▶ __half2uchar_rz (±0) returns 0.
▶ __half2uchar_rz (x), x > 255 returns UCHAR_MAX = 0xFF.
▶ __half2uchar_rz (x), x < 0.0 returns 0.
▶ __half2uchar_rz(NaN) returns 0.
__device__ unsigned int __half2uint_rd(const __half h)
Convert a half to an unsigned integer in round-down mode.
Convert the half-precision floating-point value h to an unsigned integer in round-down mode.
NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
unsigned int
▶ h converted to an unsigned integer using round-down mode.
▶ __half2uint_rd (±0) returns 0.
▶ __half2uint_rd (+∞) returns UINT_MAX = 0xFFFFFFFF.
▶ __half2uint_rd (x), x < 0.0 returns 0.
▶ __half2uint_rd(NaN) returns 0.
__device__ unsigned int __half2uint_rn(const __half h)
Convert a half to an unsigned integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to an unsigned integer in round-to-nearest-even
mode. NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
82
Chapter 4. Half Precision Intrinsics
Returns
unsigned int
▶ h converted to an unsigned integer using round-to-nearest-even mode.
▶ __half2uint_rn (±0) returns 0.
▶ __half2uint_rn (+∞) returns UINT_MAX = 0xFFFFFFFF.
▶ __half2uint_rn (x), x < 0.0 returns 0.
▶ __half2uint_rn(NaN) returns 0.
__device__ unsigned int __half2uint_ru(const __half h)
Convert a half to an unsigned integer in round-up mode.
Convert the half-precision floating-point value h to an unsigned integer in round-up mode. NaN
inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
unsigned int
▶ h converted to an unsigned integer using round-up mode.
▶ __half2uint_ru (±0) returns 0.
▶ __half2uint_ru (+∞) returns UINT_MAX = 0xFFFFFFFF.
▶ __half2uint_ru (x), x < 0.0 returns 0.
▶ __half2uint_ru(NaN) returns 0.
__host__ __device__ unsigned int __half2uint_rz(const __half h)
Convert a half to an unsigned integer in round-towards-zero mode.
Convert the half-precision floating-point value h to an unsigned integer in round-towards-zero
mode. NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
unsigned int
▶ h converted to an unsigned integer using round-towards-zero mode.
▶ __half2uint_rz (±0) returns 0.
▶ __half2uint_rz (+∞) returns UINT_MAX = 0xFFFFFFFF.
▶ __half2uint_rz (x), x < 0.0 returns 0.
▶ __half2uint_rz(NaN) returns 0.
__device__ unsigned long long int __half2ull_rd(const __half h)
Convert a half to an unsigned 64-bit integer in round-down mode.
Convert the half-precision floating-point value h to an unsigned 64-bit integer in round-down
mode. NaN inputs return 0x8000000000000000.
Parameters
h - [in] - half. Is only being read.
4.5. Half Precision Conversion and Data Movement
83
Returns
unsigned long long int
▶ h converted to an unsigned 64-bit integer using round-down mode.
▶ __half2ull_rd (±0) returns 0.
▶ __half2ull_rd (+∞) returns ULLONG_MAX = 0xFFFFFFFFFFFFFFFF.
▶ __half2ull_rd (x), x < 0.0 returns 0.
▶ __half2ull_rd(NaN) returns 0x8000000000000000.
__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.
Convert the half-precision floating-point value h to an unsigned 64-bit integer in round-to-
nearest-even mode. NaN inputs return 0x8000000000000000.
Parameters
h - [in] - half. Is only being read.
Returns
unsigned long long int
▶ h converted to an unsigned 64-bit integer using round-to-nearest-even mode.
▶ __half2ull_rn (±0) returns 0.
▶ __half2ull_rn (+∞) returns ULLONG_MAX = 0xFFFFFFFFFFFFFFFF.
▶ __half2ull_rn (x), x < 0.0 returns 0.
▶ __half2ull_rn(NaN) returns 0x8000000000000000.
__device__ unsigned long long int __half2ull_ru(const __half h)
Convert a half to an unsigned 64-bit integer in round-up mode.
Convert the half-precision floating-point value h to an unsigned 64-bit integer in round-up mode.
NaN inputs return 0x8000000000000000.
Parameters
h - [in] - half. Is only being read.
Returns
unsigned long long int
▶ h converted to an unsigned 64-bit integer using round-up mode.
▶ __half2ull_ru (±0) returns 0.
▶ __half2ull_ru (+∞) returns ULLONG_MAX = 0xFFFFFFFFFFFFFFFF.
▶ __half2ull_ru (x), x < 0.0 returns 0.
▶ __half2ull_ru(NaN) returns 0x8000000000000000.
__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.
Convert the half-precision floating-point value h to an unsigned 64-bit integer in round-towards-
zero mode. NaN inputs return 0x8000000000000000.
Parameters
h - [in] - half. Is only being read.
84
Chapter 4. Half Precision Intrinsics
Returns
unsigned long long int
▶ h converted to an unsigned 64-bit integer using round-towards-zero mode.
▶ __half2ull_rz (±0) returns 0.
▶ __half2ull_rz (+∞) returns ULLONG_MAX = 0xFFFFFFFFFFFFFFFF.
▶ __half2ull_rz (x), x < 0.0 returns 0.
▶ __half2ull_rz(NaN) returns 0x8000000000000000.
__device__ unsigned short int __half2ushort_rd(const __half h)
Convert a half to an unsigned short integer in round-down mode.
Convert the half-precision floating-point value h to an unsigned short integer in round-down
mode. NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
unsigned short int
▶ h converted to an unsigned short integer using round-down mode.
▶ __half2ushort_rd (±0) returns 0.
▶ __half2ushort_rd (+∞) returns USHRT_MAX = 0xFFFF.
▶ __half2ushort_rd (x), x < 0.0 returns 0.
▶ __half2ushort_rd(NaN) returns 0.
__device__ unsigned short int __half2ushort_rn(const __half h)
Convert a half to an unsigned short integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to an unsigned short integer in round-to-
nearest-even mode. NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
unsigned short int
▶ h converted to an unsigned short integer using round-to-nearest-even mode.
▶ __half2ushort_rn (±0) returns 0.
▶ __half2ushort_rn (+∞) returns USHRT_MAX = 0xFFFF.
▶ __half2ushort_rn (x), x < 0.0 returns 0.
▶ __half2ushort_rn(NaN) returns 0.
__device__ unsigned short int __half2ushort_ru(const __half h)
Convert a half to an unsigned short integer in round-up mode.
Convert the half-precision floating-point value h to an unsigned short integer in round-up mode.
NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
4.5. Half Precision Conversion and Data Movement
85
Returns
unsigned short int
▶ h converted to an unsigned short integer using round-up mode.
▶ __half2ushort_ru (±0) returns 0.
▶ __half2ushort_ru (+∞) returns USHRT_MAX = 0xFFFF.
▶ __half2ushort_ru (x), x < 0.0 returns 0.
▶ __half2ushort_ru(NaN) returns 0.
__host__ __device__ unsigned short int __half2ushort_rz(const __half h)
Convert a half to an unsigned short integer in round-towards-zero mode.
Convert the half-precision floating-point value h to an unsigned short integer in round-towards-
zero mode. NaN inputs are converted to 0.
Parameters
h - [in] - half. Is only being read.
Returns
unsigned short int
▶ h converted to an unsigned short integer using round-towards-zero mode.
▶ __half2ushort_rz (±0) returns 0.
▶ __half2ushort_rz (+∞) returns USHRT_MAX = 0xFFFF.
▶ __half2ushort_rz (x), x < 0.0 returns 0.
▶ __half2ushort_rz(NaN) returns 0.
__host__ __device__ short int __half_as_short(const __half h)
Reinterprets bits in a half as a signed short integer.
Reinterprets the bits in the half-precision floating-point number h as a signed short integer.
Parameters
h - [in] - half. Is only being read.
Returns
short int
▶ The reinterpreted value.
__host__ __device__ unsigned short int __half_as_ushort(const __half h)
Reinterprets bits in a half as an unsigned short integer.
Reinterprets the bits in the half-precision floating-point h as an unsigned short number.
Parameters
h - [in] - half. Is only being read.
Returns
unsigned short int
▶ The reinterpreted value.
__host__ __device__ __half2 __halves2half2(const __half a, const __half b)
Combines two half numbers into one half2 number.
Combines two input half number a and b into one half2 number. Input a is stored in low 16
bits of the return value, input b is stored in high 16 bits of the return value.
86
Chapter 4. Half Precision Intrinsics
Parameters
▶ a - [in] - half. Is only being read.
▶ b - [in] - half. Is only being read.
Returns
half2
▶ The half2 with one half equal to a and the other to b.
__host__ __device__ float __high2float(const __half2 a)
Converts high 16 bits of half2 to float and returns the result.
Converts high 16 bits of half2 input a to 32-bit floating-point number and returns the result.
See also:
__half2float(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
float
▶ The high 16 bits of a converted to float.
__host__ __device__ __half __high2half(const __half2 a)
Returns high 16 bits of half2 input.
Returns high 16 bits of half2 input a.
Parameters
a - [in] - half2. Is only being read.
Returns
half
▶ The high 16 bits of the input.
__host__ __device__ __half2 __high2half2(const __half2 a)
Extracts high 16 bits from half2 input.
Extracts high 16 bits from half2 input a and returns a new half2 number which has both halves
equal to the extracted bits.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The half2 with both halves equal to the high 16 bits of the 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.
Extracts high 16 bits from each of the two half2 inputs and combines into one half2 number.
High 16 bits from input a is stored in low 16 bits of the return value, high 16 bits from input b is
stored in high 16 bits of the return value.
4.5. Half Precision Conversion and Data Movement
87
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The high 16 bits of a and of b.
__host__ __device__ __half __int2half_rd(const int i)
Convert a signed integer to a half in round-down mode.
Convert the signed integer value i to a half-precision floating-point value in round-down mode.
Parameters
i - [in] - int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __int2half_rn(const int i)
Convert a signed integer to a half in round-to-nearest-even mode.
Convert the signed integer value i to a half-precision floating-point value in round-to-nearest-
even mode.
Parameters
i - [in] - int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __int2half_ru(const int i)
Convert a signed integer to a half in round-up mode.
Convert the signed integer value i to a half-precision floating-point value in round-up mode.
Parameters
i - [in] - int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __int2half_rz(const int i)
Convert a signed integer to a half in round-towards-zero mode.
Convert the signed integer value i to a half-precision floating-point value in round-towards-zero
mode.
Parameters
i - [in] - int. Is only being read.
Returns
half
▶ i converted to half.
88
Chapter 4. Half Precision Intrinsics
__device__ __half2 __ldca(const __half2 *const ptr)
Generates a ld.global.ca load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __half __ldca(const __half *const ptr)
Generates a ld.global.ca load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __half __ldcg(const __half *const ptr)
Generates a ld.global.cg load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __half2 __ldcg(const __half2 *const ptr)
Generates a ld.global.cg load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __half __ldcs(const __half *const ptr)
Generates a ld.global.cs load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __half2 __ldcs(const __half2 *const ptr)
Generates a ld.global.cs load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __half2 __ldcv(const __half2 *const ptr)
Generates a ld.global.cv load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
4.5. Half Precision Conversion and Data Movement
89
__device__ __half __ldcv(const __half *const ptr)
Generates a ld.global.cv load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __half2 __ldg(const __half2 *const ptr)
Generates a ld.global.nc load instruction.
defined(CUDA_ARCH) || (CUDA_ARCH >= 300)
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __half __ldg(const __half *const ptr)
Generates a ld.global.nc load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __half __ldlu(const __half *const ptr)
Generates a ld.global.lu load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __half2 __ldlu(const __half2 *const ptr)
Generates a ld.global.lu load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__host__ __device__ __half __ll2half_rd(const long long int i)
Convert a signed 64-bit integer to a half in round-down mode.
Convert the signed 64-bit integer value i to a half-precision floating-point value in round-down
mode.
Parameters
i - [in] - long long int. Is only being read.
Returns
half
▶ i converted to half.
90
Chapter 4. Half Precision Intrinsics
__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.
Convert the signed 64-bit integer value i to a half-precision floating-point value in round-to-
nearest-even mode.
Parameters
i - [in] - long long int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __ll2half_ru(const long long int i)
Convert a signed 64-bit integer to a half in round-up mode.
Convert the signed 64-bit integer value i to a half-precision floating-point value in round-up
mode.
Parameters
i - [in] - long long int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __ll2half_rz(const long long int i)
Convert a signed 64-bit integer to a half in round-towards-zero mode.
Convert the signed 64-bit integer value i to a half-precision floating-point value in round-
towards-zero mode.
Parameters
i - [in] - long long int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ float __low2float(const __half2 a)
Converts low 16 bits of half2 to float and returns the result.
Converts low 16 bits of half2 input a to 32-bit floating-point number and returns the result.
See also:
__half2float(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
float
▶ The low 16 bits of a converted to float.
__host__ __device__ __half __low2half(const __half2 a)
Returns low 16 bits of half2 input.
Returns low 16 bits of half2 input a.
4.5. Half Precision Conversion and Data Movement
91
Parameters
a - [in] - half2. Is only being read.
Returns
half
▶ Returns half which contains low 16 bits of the input a.
__host__ __device__ __half2 __low2half2(const __half2 a)
Extracts low 16 bits from half2 input.
Extracts low 16 bits from half2 input a and returns a new half2 number which has both halves
equal to the extracted bits.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The half2 with both halves equal to the low 16 bits of the input.
__host__ __device__ __half2 __lowhigh2highlow(const __half2 a)
Swaps both halves of the half2 input.
Swaps both halves of the half2 input and returns a new half2 number with swapped halves.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ a with its halves being swapped.
__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.
Extracts low 16 bits from each of the two half2 inputs and combines into one half2 number.
Low 16 bits from input a is stored in low 16 bits of the return value, low 16 bits from input b is
stored in high 16 bits of the return value.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The low 16 bits of a and of b.
__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.
Copy from a thread with higher ID relative to the caller.
Calculates a source thread ID by adding delta to the caller’s thread ID. The value of var held by
the resulting thread ID is returned: this has the effect of shifting var down the warp by delta
threads. If the width is less than warpSize, then each subsection of the warp behaves as a sep-
arate entity with a starting logical thread ID of 0. Similarly to the __shfl_up_sync(), the ID number
92
Chapter 4. Half Precision Intrinsics
of the source thread will not wrap around the value of width and the upper delta threads will
remain unchanged. Threads may only read data from another thread which is actively partici-
pating in the __shfl_*sync() command. If the target thread is inactive, the retrieved value is
undefined.
Note: For more details about this function, see the Warp Shuffle Functions section in the CUDA
C++ Programming Guide.
Parameters
▶ mask - [in] - unsigned int. Is only being read.
▶ Indicates the threads participating in the call.
▶ A bit, representing the thread’s lane ID, must be set for each participating
thread to ensure they are properly converged before the intrinsic is executed
by the hardware.
▶ Each calling thread must have its own bit set in the mask and all non-exited
threads named in mask must execute the same intrinsic with the same mask,
or the result is undefined.
▶ var - [in] - half. Is only being read.
▶ delta - [in] - unsigned int. Is only being read.
▶ width - [in] - int. Is only being read.
Returns
Returns the 2-byte word referenced by var from the source thread ID as half.
__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.
Copy from a thread with higher ID relative to the caller.
Calculates a source thread ID by adding delta to the caller’s thread ID. The value of var held by
the resulting thread ID is returned: this has the effect of shifting var down the warp by delta
threads. If the width is less than warpSize, then each subsection of the warp behaves as a sep-
arate entity with a starting logical thread ID of 0. Similarly to the __shfl_up_sync(), the ID number
of the source thread will not wrap around the value of width and the upper delta threads will
remain unchanged. Threads may only read data from another thread which is actively partici-
pating in the __shfl_*sync() command. If the target thread is inactive, the retrieved value is
undefined.
Note: For more details about this function, see the Warp Shuffle Functions section in the CUDA
C++ Programming Guide.
Parameters
▶ mask - [in] - unsigned int. Is only being read.
▶ Indicates the threads participating in the call.
4.5. Half Precision Conversion and Data Movement
93
▶ A bit, representing the thread’s lane ID, must be set for each participating
thread to ensure they are properly converged before the intrinsic is executed
by the hardware.
▶ Each calling thread must have its own bit set in the mask and all non-exited
threads named in mask must execute the same intrinsic with the same mask,
or the result is undefined.
▶ var - [in] - half2. Is only being read.
▶ delta - [in] - unsigned int. Is only being read.
▶ width - [in] - int. Is only being read.
Returns
Returns the 4-byte word referenced by var from the source thread ID as half2.
__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.
Direct copy from indexed thread.
Returns the value of var held by the thread whose ID is given by srcLane. If the width is less
than warpSize, then each subsection of the warp behaves as a separate entity with a starting
logical thread ID of 0. If srcLane is outside the range [0:width-1], the value returned corresponds
to the value of var held by the srcLane modulo width (i.e. within the same subsection). width
must have a value which is a power of 2; results are undefined if width is not a power of 2, or
is a number greater than warpSize. Threads may only read data from another thread which
is actively participating in the __shfl_*sync() command. If the target thread is inactive, the
retrieved value is undefined.
Note: For more details about this function, see the Warp Shuffle Functions section in the CUDA
C++ Programming Guide.
Parameters
▶ mask - [in] - unsigned int. Is only being read.
▶ Indicates the threads participating in the call.
▶ A bit, representing the thread’s lane ID, must be set for each participating
thread to ensure they are properly converged before the intrinsic is executed
by the hardware.
▶ Each calling thread must have its own bit set in the mask and all non-exited
threads named in mask must execute the same intrinsic with the same mask,
or the result is undefined.
▶ var - [in] - half2. Is only being read.
▶ srcLane - [in] - int. Is only being read.
▶ width - [in] - int. Is only being read.
Returns
Returns the 4-byte word referenced by var from the source thread ID as half2.
94
Chapter 4. Half Precision Intrinsics
__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.
Direct copy from indexed thread.
Returns the value of var held by the thread whose ID is given by srcLane. If the width is less
than warpSize, then each subsection of the warp behaves as a separate entity with a starting
logical thread ID of 0. If srcLane is outside the range [0:width-1], the value returned corresponds
to the value of var held by the srcLane modulo width (i.e. within the same subsection). width
must have a value which is a power of 2; results are undefined if width is not a power of 2, or
is a number greater than warpSize. Threads may only read data from another thread which
is actively participating in the __shfl_*sync() command. If the target thread is inactive, the
retrieved value is undefined.
Note: For more details about this function, see the Warp Shuffle Functions section in the CUDA
C++ Programming Guide.
Parameters
▶ mask - [in] - unsigned int. Is only being read.
▶ Indicates the threads participating in the call.
▶ A bit, representing the thread’s lane ID, must be set for each participating
thread to ensure they are properly converged before the intrinsic is executed
by the hardware.
▶ Each calling thread must have its own bit set in the mask and all non-exited
threads named in mask must execute the same intrinsic with the same mask,
or the result is undefined.
▶ var - [in] - half. Is only being read.
▶ srcLane - [in] - int. Is only being read.
▶ width - [in] - int. Is only being read.
Returns
Returns the 2-byte word referenced by var from the source thread ID as half.
__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.
Copy from a thread with lower ID relative to the caller.
Calculates a source thread ID by subtracting delta from the caller’s lane ID. The value of var
held by the resulting lane ID is returned: in effect, var is shifted up the warp by delta threads. If
the width is less than warpSize, then each subsection of the warp behaves as a separate entity
with a starting logical thread ID of 0. The source thread index will not wrap around the value of
width, so effectively the lower delta threads will be unchanged. width must have a value which
is a power of 2; results are undefined if width is not a power of 2, or is a number greater than
warpSize. Threads may only read data from another thread which is actively participating in the
__shfl_*sync() command. If the target thread is inactive, the retrieved value is undefined.
4.5. Half Precision Conversion and Data Movement
95
Note: For more details about this function, see the Warp Shuffle Functions section in the CUDA
C++ Programming Guide.
Parameters
▶ mask - [in] - unsigned int. Is only being read.
▶ Indicates the threads participating in the call.
▶ A bit, representing the thread’s lane ID, must be set for each participating
thread to ensure they are properly converged before the intrinsic is executed
by the hardware.
▶ Each calling thread must have its own bit set in the mask and all non-exited
threads named in mask must execute the same intrinsic with the same mask,
or the result is undefined.
▶ var - [in] - half2. Is only being read.
▶ delta - [in] - unsigned int. Is only being read.
▶ width - [in] - int. Is only being read.
Returns
Returns the 4-byte word referenced by var from the source thread ID as half2.
__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.
Copy from a thread with lower ID relative to the caller.
Calculates a source thread ID by subtracting delta from the caller’s lane ID. The value of var
held by the resulting lane ID is returned: in effect, var is shifted up the warp by delta threads. If
the width is less than warpSize, then each subsection of the warp behaves as a separate entity
with a starting logical thread ID of 0. The source thread index will not wrap around the value of
width, so effectively the lower delta threads will be unchanged. width must have a value which
is a power of 2; results are undefined if width is not a power of 2, or is a number greater than
warpSize. Threads may only read data from another thread which is actively participating in the
__shfl_*sync() command. If the target thread is inactive, the retrieved value is undefined.
Note: For more details about this function, see the Warp Shuffle Functions section in the CUDA
C++ Programming Guide.
Parameters
▶ mask - [in] - unsigned int. Is only being read.
▶ Indicates the threads participating in the call.
▶ A bit, representing the thread’s lane ID, must be set for each participating
thread to ensure they are properly converged before the intrinsic is executed
by the hardware.
96
Chapter 4. Half Precision Intrinsics
▶ Each calling thread must have its own bit set in the mask and all non-exited
threads named in mask must execute the same intrinsic with the same mask,
or the result is undefined.
▶ var - [in] - half. Is only being read.
▶ delta - [in] - unsigned int. Is only being read.
▶ width - [in] - int. Is only being read.
Returns
Returns the 2-byte word referenced by var from the source thread ID as half.
__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.
Copy from a thread based on bitwise XOR of own thread ID.
Calculates a source thread ID by performing a bitwise XOR of the caller’s thread ID with lane-
Mask: the value of var held by the resulting thread ID is returned. If the width is less than
warpSize, then each group of width consecutive threads are able to access elements from ear-
lier groups of threads, however if they attempt to access elements from later groups of threads
their own value of var will be returned. This mode implements a butterfly addressing pattern
such as is used in tree reduction and broadcast. Threads may only read data from another thread
which is actively participating in the __shfl_*sync() command. If the target thread is inactive,
the retrieved value is undefined.
Note: For more details about this function, see the Warp Shuffle Functions section in the CUDA
C++ Programming Guide.
Parameters
▶ mask - [in] - unsigned int. Is only being read.
▶ Indicates the threads participating in the call.
▶ A bit, representing the thread’s lane ID, must be set for each participating
thread to ensure they are properly converged before the intrinsic is executed
by the hardware.
▶ Each calling thread must have its own bit set in the mask and all non-exited
threads named in mask must execute the same intrinsic with the same mask,
or the result is undefined.
▶ var - [in] - half2. Is only being read.
▶ laneMask - [in] - int. Is only being read.
▶ width - [in] - int. Is only being read.
Returns
Returns the 4-byte word referenced by var from the source thread ID as half2.
__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.
4.5. Half Precision Conversion and Data Movement
97
Copy from a thread based on bitwise XOR of own thread ID.
Calculates a source thread ID by performing a bitwise XOR of the caller’s thread ID with lane-
Mask: the value of var held by the resulting thread ID is returned. If the width is less than
warpSize, then each group of width consecutive threads are able to access elements from ear-
lier groups of threads, however if they attempt to access elements from later groups of threads
their own value of var will be returned. This mode implements a butterfly addressing pattern
such as is used in tree reduction and broadcast. Threads may only read data from another thread
which is actively participating in the __shfl_*sync() command. If the target thread is inactive,
the retrieved value is undefined.
Note: For more details about this function, see the Warp Shuffle Functions section in the CUDA
C++ Programming Guide.
Parameters
▶ mask - [in] - unsigned int. Is only being read.
▶ Indicates the threads participating in the call.
▶ A bit, representing the thread’s lane ID, must be set for each participating
thread to ensure they are properly converged before the intrinsic is executed
by the hardware.
▶ Each calling thread must have its own bit set in the mask and all non-exited
threads named in mask must execute the same intrinsic with the same mask,
or the result is undefined.
▶ var - [in] - half. Is only being read.
▶ laneMask - [in] - int. Is only being read.
▶ width - [in] - int. Is only being read.
Returns
Returns the 2-byte word referenced by var from the source thread ID as half.
__host__ __device__ __half __short2half_rd(const short int i)
Convert a signed short integer to a half in round-down mode.
Convert the signed short integer value i to a half-precision floating-point value in round-down
mode.
Parameters
i - [in] - short int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __short2half_rn(const short int i)
Convert a signed short integer to a half in round-to-nearest-even mode.
Convert the signed short integer value i to a half-precision floating-point value in round-to-
nearest-even mode.
Parameters
i - [in] - short int. Is only being read.
98
Chapter 4. Half Precision Intrinsics
Returns
half
▶ i converted to half.
__host__ __device__ __half __short2half_ru(const short int i)
Convert a signed short integer to a half in round-up mode.
Convert the signed short integer value i to a half-precision floating-point value in round-up mode.
Parameters
i - [in] - short int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __short2half_rz(const short int i)
Convert a signed short integer to a half in round-towards-zero mode.
Convert the signed short integer value i to a half-precision floating-point value in round-towards-
zero mode.
Parameters
i - [in] - short int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __short_as_half(const short int i)
Reinterprets bits in a signed short integer as a half.
Reinterprets the bits in the signed short integer i as a half-precision floating-point number.
Parameters
i - [in] - short int. Is only being read.
Returns
half
▶ The reinterpreted value.
__device__ void __stcg(__half2 *const ptr, const __half2 value)
Generates a st.global.cg store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stcg(__half *const ptr, const __half value)
Generates a st.global.cg store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
4.5. Half Precision Conversion and Data Movement
99
__device__ void __stcs(__half2 *const ptr, const __half2 value)
Generates a st.global.cs store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stcs(__half *const ptr, const __half value)
Generates a st.global.cs store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stwb(__half2 *const ptr, const __half2 value)
Generates a st.global.wb store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stwb(__half *const ptr, const __half value)
Generates a st.global.wb store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stwt(__half *const ptr, const __half value)
Generates a st.global.wt store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stwt(__half2 *const ptr, const __half2 value)
Generates a st.global.wt store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__host__ __device__ __half __uint2half_rd(const unsigned int i)
Convert an unsigned integer to a half in round-down mode.
Convert the unsigned integer value i to a half-precision floating-point value in round-down mode.
100
Chapter 4. Half Precision Intrinsics
Parameters
i - [in] - unsigned int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __uint2half_rn(const unsigned int i)
Convert an unsigned integer to a half in round-to-nearest-even mode.
Convert the unsigned integer value i to a half-precision floating-point value in round-to-nearest-
even mode.
Parameters
i - [in] - unsigned int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __uint2half_ru(const unsigned int i)
Convert an unsigned integer to a half in round-up mode.
Convert the unsigned integer value i to a half-precision floating-point value in round-up mode.
Parameters
i - [in] - unsigned int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __uint2half_rz(const unsigned int i)
Convert an unsigned integer to a half in round-towards-zero mode.
Convert the unsigned integer value i to a half-precision floating-point value in round-towards-
zero mode.
Parameters
i - [in] - unsigned int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __ull2half_rd(const unsigned long long int i)
Convert an unsigned 64-bit integer to a half in round-down mode.
Convert the unsigned 64-bit integer value i to a half-precision floating-point value in round-down
mode.
Parameters
i - [in] - unsigned long long int. Is only being read.
Returns
half
▶ i converted to half.
4.5. Half Precision Conversion and Data Movement
101
__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.
Convert the unsigned 64-bit integer value i to a half-precision floating-point value in round-to-
nearest-even mode.
Parameters
i - [in] - unsigned long long int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __ull2half_ru(const unsigned long long int i)
Convert an unsigned 64-bit integer to a half in round-up mode.
Convert the unsigned 64-bit integer value i to a half-precision floating-point value in round-up
mode.
Parameters
i - [in] - unsigned long long int. Is only being read.
Returns
half
▶ i converted to half.
__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.
Convert the unsigned 64-bit integer value i to a half-precision floating-point value in round-
towards-zero mode.
Parameters
i - [in] - unsigned long long int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __ushort2half_rd(const unsigned short int i)
Convert an unsigned short integer to a half in round-down mode.
Convert the unsigned short integer value i to a half-precision floating-point value in round-down
mode.
Parameters
i - [in] - unsigned short int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __ushort2half_rn(const unsigned short int i)
Convert an unsigned short integer to a half in round-to-nearest-even mode.
Convert the unsigned short integer value i to a half-precision floating-point value in round-to-
nearest-even mode.
Parameters
i - [in] - unsigned short int. Is only being read.
102
Chapter 4. Half Precision Intrinsics
Returns
half
▶ i converted to half.
__host__ __device__ __half __ushort2half_ru(const unsigned short int i)
Convert an unsigned short integer to a half in round-up mode.
Convert the unsigned short integer value i to a half-precision floating-point value in round-up
mode.
Parameters
i - [in] - unsigned short int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __ushort2half_rz(const unsigned short int i)
Convert an unsigned short integer to a half in round-towards-zero mode.
Convert the unsigned short integer value i to a half-precision floating-point value in round-
towards-zero mode.
Parameters
i - [in] - unsigned short int. Is only being read.
Returns
half
▶ i converted to half.
__host__ __device__ __half __ushort_as_half(const unsigned short int i)
Reinterprets bits in an unsigned short integer as a half.
Reinterprets the bits in the unsigned short integer i as a half-precision floating-point number.
Parameters
i - [in] - unsigned short int. Is only being read.
Returns
half
▶ The reinterpreted value.
__host__ __device__ __half2 make_half2(const __half x, const __half y)
Vector function, combines two __half numbers into one __half2 number.
Combines two input __half number x and y into one __half2 number. Input x is stored in low
16 bits of the return value, input y is stored in high 16 bits of the return value.
Parameters
▶ x - [in] - half. Is only being read.
▶ y - [in] - half. Is only being read.
Returns
__half2
▶ The __half2 vector with one half equal to x and the other to y.
4.5. Half Precision Conversion and Data Movement
103
4.6. Half2 Arithmetic Functions
To use these functions, include the header file cuda_fp16.h in your program.
Functions
__host__ __device__ __half2 __h2div(const __half2 a, const __half2 b)
Performs half2 vector division in round-to-nearest-even mode.
__host__ __device__ __half2 __habs2(const __half2 a)
Calculates the absolute value of both halves of the input half2 number and returns the result.
__host__ __device__ __half2 __hadd2(const __half2 a, const __half2 b)
Performs half2 vector addition in round-to-nearest-even mode.
__host__ __device__ __half2 __hadd2_rn(const __half2 a, const __half2 b)
Performs half2 vector addition in round-to-nearest-even mode.
__host__ __device__ __half2 __hadd2_sat(const __half2 a, const __half2 b)
Performs half2 vector addition in round-to-nearest-even mode, with saturation to [0.0, 1.0].
__device__ __half2 __hcmadd(const __half2 a, const __half2 b, const __half2 c)
Performs fast complex multiply-accumulate.
__device__ __half2 __hfma2(const __half2 a, const __half2 b, const __half2 c)
Performs half2 vector fused multiply-add in round-to-nearest-even mode.
__device__ __half2 __hfma2_relu(const __half2 a, const __half2 b, const __half2 c)
Performs half2 vector fused multiply-add in round-to-nearest-even mode with relu saturation.
__device__ __half2 __hfma2_sat(const __half2 a, const __half2 b, const __half2 c)
Performs half2 vector fused multiply-add in round-to-nearest-even mode, with saturation to
[0.0, 1.0].
__host__ __device__ __half2 __hmul2(const __half2 a, const __half2 b)
Performs half2 vector multiplication in round-to-nearest-even mode.
__host__ __device__ __half2 __hmul2_rn(const __half2 a, const __half2 b)
Performs half2 vector multiplication in round-to-nearest-even mode.
__host__ __device__ __half2 __hmul2_sat(const __half2 a, const __half2 b)
Performs half2 vector multiplication in round-to-nearest-even mode, with saturation to [0.0,
1.0].
__host__ __device__ __half2 __hneg2(const __half2 a)
Negates both halves of the input half2 number and returns the result.
__host__ __device__ __half2 __hsub2(const __half2 a, const __half2 b)
Performs half2 vector subtraction in round-to-nearest-even mode.
__host__ __device__ __half2 __hsub2_rn(const __half2 a, const __half2 b)
Performs half2 vector subtraction in round-to-nearest-even mode.
__host__ __device__ __half2 __hsub2_sat(const __half2 a, const __half2 b)
Performs half2 vector subtraction in round-to-nearest-even mode, with saturation to [0.0, 1.0].
__device__ __half2 atomicAdd(__half2 *const address, const __half2 val)
Vector add val to the value stored at address in global or shared memory, and writes this value
back to address .
104
Chapter 4. Half Precision Intrinsics
__host__ __device__ __half2 operator*(const __half2 &lh, const __half2 &rh)
Performs packed half multiplication operation.
__host__ __device__ __half2 & operator*=(__half2 &lh, const __half2 &rh)
Performs packed half compound assignment with multiplication operation.
__host__ __device__ __half2 operator+(const __half2 &h)
Implements packed half unary plus operator, returns input value.
__host__ __device__ __half2 operator+(const __half2 &lh, const __half2 &rh)
Performs packed half addition operation.
__host__ __device__ __half2 operator++(__half2 &h, const int ignored)
Performs packed half postfix increment operation.
__host__ __device__ __half2 & operator++(__half2 &h)
Performs packed half prefix increment operation.
__host__ __device__ __half2 & operator+=(__half2 &lh, const __half2 &rh)
Performs packed half compound assignment with addition operation.
__host__ __device__ __half2 operator-(const __half2 &h)
Implements packed half unary minus operator.
__host__ __device__ __half2 operator-(const __half2 &lh, const __half2 &rh)
Performs packed half subtraction operation.
__host__ __device__ __half2 & operator-(__half2 &h)
Performs packed half prefix decrement operation.
__host__ __device__ __half2 operator-(__half2 &h, const int ignored)
Performs packed half postfix decrement operation.
__host__ __device__ __half2 & operator-=(__half2 &lh, const __half2 &rh)
Performs packed half compound assignment with subtraction operation.
__host__ __device__ __half2 operator/(const __half2 &lh, const __half2 &rh)
Performs packed half division operation.
__host__ __device__ __half2 & operator/=(__half2 &lh, const __half2 &rh)
Performs packed half compound assignment with division operation.
4.6.1. Functions
__host__ __device__ __half2 __h2div(const __half2 a, const __half2 b)
Performs half2 vector division in round-to-nearest-even mode.
Divides half2 input vector a by input vector b in round-to-nearest-even mode.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise division of a with b.
4.6. Half2 Arithmetic Functions
105
__host__ __device__ __half2 __habs2(const __half2 a)
Calculates the absolute value of both halves of the input half2 number and returns the result.
Calculates the absolute value of both halves of the input half2 number and returns the result.
See also:
__habs(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ Returns a with the absolute value of both halves.
__host__ __device__ __half2 __hadd2(const __half2 a, const __half2 b)
Performs half2 vector addition in round-to-nearest-even mode.
Performs half2 vector add of inputs a and b, in round-to-nearest-even mode.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The sum of vectors a and b.
__host__ __device__ __half2 __hadd2_rn(const __half2 a, const __half2 b)
Performs half2 vector addition in round-to-nearest-even mode.
Performs half2 vector add of inputs a and b, in round-to-nearest-even mode. Prevents floating-
point contractions of mul+add into fma.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The sum of vectors a and b.
__host__ __device__ __half2 __hadd2_sat(const __half2 a, const __half2 b)
Performs half2 vector addition in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs half2 vector add of inputs a and b, in round-to-nearest-even mode, and clamps the
results to range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
106
Chapter 4. Half Precision Intrinsics
Returns
half2
▶ The sum of a and b, with respect to saturation.
__device__ __half2 __hcmadd(const __half2 a, const __half2 b, const __half2 c)
Performs fast complex multiply-accumulate.
Interprets vector half2 input pairs a, b, and c as complex numbers in half precision: (a.x + I*a.y),
(b.x + I*b.y), (c.x + I*c.y) and performs complex multiply-accumulate operation: a*b + c in a simple
way: ((a.x*b.x + c.x) - a.y*b.y) + I*((a.x*b.y + c.y) + a.y*b.x)
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
▶ c - [in] - half2. Is only being read.
Returns
half2
▶ The result of complex multiply-accumulate operation on complex numbers a, b,
and c
▶ __half2 result = __hcmadd(a, b, c) is numerically in agreement with:
▶ result.x = __hfma(-a.y, b.y, __hfma(a.x, b.x, c.x))
▶ result.y = __hfma( a.y, b.x, __hfma(a.x, b.y, c.y))
__device__ __half2 __hfma2(const __half2 a, const __half2 b, const __half2 c)
Performs half2 vector fused multiply-add in round-to-nearest-even mode.
Performs half2 vector multiply on inputs a and b, then performs a half2 vector add of the
result with c, rounding the result once in round-to-nearest-even mode.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
▶ c - [in] - half2. Is only being read.
Returns
half2
▶ The result of elementwise fused multiply-add operation on vectors a, b, and c.
__device__ __half2 __hfma2_relu(const __half2 a, const __half2 b, const __half2 c)
Performs half2 vector fused multiply-add in round-to-nearest-even mode with relu saturation.
Performs half2 vector multiply on inputs a and b, then performs a half2 vector 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] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
4.6. Half2 Arithmetic Functions
107
▶ c - [in] - half2. Is only being read.
Returns
half2
▶ The result of elementwise fused multiply-add operation on vectors a, b, and c
with relu saturation.
__device__ __half2 __hfma2_sat(const __half2 a, const __half2 b, const __half2 c)
Performs half2 vector fused multiply-add in round-to-nearest-even mode, with saturation to
[0.0, 1.0].
Performs half2 vector multiply on inputs a and b, then performs a half2 vector add of the
result with c, rounding the result once in round-to-nearest-even mode, and clamps the results to
range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
▶ c - [in] - half2. Is only being read.
Returns
half2
▶ The result of elementwise fused multiply-add operation on vectors a, b, and c,
with respect to saturation.
__host__ __device__ __half2 __hmul2(const __half2 a, const __half2 b)
Performs half2 vector multiplication in round-to-nearest-even mode.
Performs half2 vector multiplication of inputs a and b, in round-to-nearest-even mode.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The result of elementwise multiplying the vectors a and b.
__host__ __device__ __half2 __hmul2_rn(const __half2 a, const __half2 b)
Performs half2 vector multiplication in round-to-nearest-even mode.
Performs half2 vector 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] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The result of elementwise multiplying the vectors a and b.
108
Chapter 4. Half Precision Intrinsics
__host__ __device__ __half2 __hmul2_sat(const __half2 a, const __half2 b)
Performs half2 vector multiplication in round-to-nearest-even mode, with saturation to [0.0,
1.0].
Performs half2 vector multiplication of inputs a and b, in round-to-nearest-even mode, and
clamps the results to range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The result of elementwise multiplication of vectors a and b, with respect to
saturation.
__host__ __device__ __half2 __hneg2(const __half2 a)
Negates both halves of the input half2 number and returns the result.
Negates both halves of the input half2 number a and returns the result.
See also:
__hneg(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ Returns a with both halves negated.
__host__ __device__ __half2 __hsub2(const __half2 a, const __half2 b)
Performs half2 vector subtraction in round-to-nearest-even mode.
Subtracts half2 input vector b from input vector a in round-to-nearest-even mode.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The subtraction of vector b from a.
__host__ __device__ __half2 __hsub2_rn(const __half2 a, const __half2 b)
Performs half2 vector subtraction in round-to-nearest-even mode.
Subtracts half2 input vector b from input vector a in round-to-nearest-even mode. Prevents
floating-point contractions of mul+sub into fma.
Parameters
4.6. Half2 Arithmetic Functions
109
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The subtraction of vector b from a.
__host__ __device__ __half2 __hsub2_sat(const __half2 a, const __half2 b)
Performs half2 vector subtraction in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Subtracts half2 input vector b from input vector a in round-to-nearest-even mode, and clamps
the results to range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The subtraction of vector b from a, with respect to saturation.
__device__ __half2 atomicAdd(__half2 *const address, const __half2 val)
Vector add val to the value stored at address in global or shared memory, and writes this value
back to address.
The atomicity of the add operation is guaranteed separately for each of the two __half elements;
the entire __half2 is not guaranteed to be atomic as a single 32-bit access.
The location of address must be in global or shared memory. This operation has undefined
behavior otherwise. This operation is natively supported by devices of compute capability 6.x
and higher, older devices use emulation path.
Note: For more details about this function, see the Atomic Functions section in the CUDA C++
Programming Guide.
Parameters
▶ address - [in] - half2*. An address in global or shared memory.
▶ val - [in] - half2. The value to be added.
Returns
half2
▶ The old value read from address.
__host__ __device__ __half2 operator*(const __half2 &lh, const __half2 &rh)
Performs packed half multiplication operation.
See also:
__hmul2(__half2, __half2)
110
Chapter 4. Half Precision Intrinsics
__host__ __device__ __half2 &operator*=(__half2 &lh, const __half2 &rh)
Performs packed half compound assignment with multiplication operation.
See also:
__hmul2(__half2, __half2)
__host__ __device__ __half2 operator+(const __half2 &h)
Implements packed half unary plus operator, returns input value.
__host__ __device__ __half2 operator+(const __half2 &lh, const __half2 &rh)
Performs packed half addition operation.
See also:
__hadd2(__half2, __half2)
__host__ __device__ __half2 operator++(__half2 &h, const int ignored)
Performs packed half postfix increment operation.
See also:
__hadd2(__half2, __half2)
__host__ __device__ __half2 &operator++(__half2 &h)
Performs packed half prefix increment operation.
See also:
__hadd2(__half2, __half2)
__host__ __device__ __half2 &operator+=(__half2 &lh, const __half2 &rh)
Performs packed half compound assignment with addition operation.
See also:
__hadd2(__half2, __half2)
__host__ __device__ __half2 operator-(const __half2 &h)
Implements packed half unary minus operator.
See also:
__hneg2(__half2)
__host__ __device__ __half2 operator-(const __half2 &lh, const __half2 &rh)
Performs packed half subtraction operation.
See also:
__hsub2(__half2, __half2)
4.6. Half2 Arithmetic Functions
111
__host__ __device__ __half2 &operator--(__half2 &h)
Performs packed half prefix decrement operation.
See also:
__hsub2(__half2, __half2)
__host__ __device__ __half2 operator--(__half2 &h, const int ignored)
Performs packed half postfix decrement operation.
See also:
__hsub2(__half2, __half2)
__host__ __device__ __half2 &operator-=(__half2 &lh, const __half2 &rh)
Performs packed half compound assignment with subtraction operation.
See also:
__hsub2(__half2, __half2)
__host__ __device__ __half2 operator∕(const __half2 &lh, const __half2 &rh)
Performs packed half division operation.
See also:
__h2div(__half2, __half2)
__host__ __device__ __half2 &operator∕=(__half2 &lh, const __half2 &rh)
Performs packed half compound assignment with division operation.
See also:
__h2div(__half2, __half2)
4.7. Half2 Comparison Functions
To use these functions, include the header file cuda_fp16.h in your program.
112
Chapter 4. Half Precision Intrinsics
Functions
__host__ __device__ bool __hbeq2(const __half2 a, const __half2 b)
Performs half2 vector if-equal comparison and returns boolean true if both half results are
true, boolean false otherwise.
__host__ __device__ bool __hbequ2(const __half2 a, const __half2 b)
Performs half2 vector unordered if-equal comparison and returns boolean true if both half
results are true, boolean false otherwise.
__host__ __device__ bool __hbge2(const __half2 a, const __half2 b)
Performs half2 vector greater-equal comparison and returns boolean true if both half results
are true, boolean false otherwise.
__host__ __device__ bool __hbgeu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-equal comparison and returns boolean true if both
half results are true, boolean false otherwise.
__host__ __device__ bool __hbgt2(const __half2 a, const __half2 b)
Performs half2 vector greater-than comparison and returns boolean true if both half results
are true, boolean false otherwise.
__host__ __device__ bool __hbgtu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-than comparison and returns boolean true if both
half results are true, boolean false otherwise.
__host__ __device__ bool __hble2(const __half2 a, const __half2 b)
Performs half2 vector less-equal comparison and returns boolean true if both half results are
true, boolean false otherwise.
__host__ __device__ bool __hbleu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-equal comparison and returns boolean true if both half
results are true, boolean false otherwise.
__host__ __device__ bool __hblt2(const __half2 a, const __half2 b)
Performs half2 vector less-than comparison and returns boolean true if both half results are
true, boolean false otherwise.
__host__ __device__ bool __hbltu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-than comparison and returns boolean true if both half
results are true, boolean false otherwise.
__host__ __device__ bool __hbne2(const __half2 a, const __half2 b)
Performs half2 vector not-equal comparison and returns boolean true if both half results are
true, boolean false otherwise.
__host__ __device__ bool __hbneu2(const __half2 a, const __half2 b)
Performs half2 vector unordered not-equal comparison and returns boolean true if both half
results are true, boolean false otherwise.
__host__ __device__ __half2 __heq2(const __half2 a, const __half2 b)
Performs half2 vector if-equal comparison.
__host__ __device__ unsigned int __heq2_mask(const __half2 a, const __half2 b)
Performs half2 vector if-equal comparison.
__host__ __device__ __half2 __hequ2(const __half2 a, const __half2 b)
Performs half2 vector unordered if-equal comparison.
__host__ __device__ unsigned int __hequ2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered if-equal comparison.
4.7. Half2 Comparison Functions
113
__host__ __device__ __half2 __hge2(const __half2 a, const __half2 b)
Performs half2 vector greater-equal comparison.
__host__ __device__ unsigned int __hge2_mask(const __half2 a, const __half2 b)
Performs half2 vector greater-equal comparison.
__host__ __device__ __half2 __hgeu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-equal comparison.
__host__ __device__ unsigned int __hgeu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-equal comparison.
__host__ __device__ __half2 __hgt2(const __half2 a, const __half2 b)
Performs half2 vector greater-than comparison.
__host__ __device__ unsigned int __hgt2_mask(const __half2 a, const __half2 b)
Performs half2 vector greater-than comparison.
__host__ __device__ __half2 __hgtu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-than comparison.
__host__ __device__ unsigned int __hgtu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-than comparison.
__host__ __device__ __half2 __hisnan2(const __half2 a)
Determine whether half2 argument is a NaN.
__host__ __device__ __half2 __hle2(const __half2 a, const __half2 b)
Performs half2 vector less-equal comparison.
__host__ __device__ unsigned int __hle2_mask(const __half2 a, const __half2 b)
Performs half2 vector less-equal comparison.
__host__ __device__ __half2 __hleu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-equal comparison.
__host__ __device__ unsigned int __hleu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered less-equal comparison.
__host__ __device__ __half2 __hlt2(const __half2 a, const __half2 b)
Performs half2 vector less-than comparison.
__host__ __device__ unsigned int __hlt2_mask(const __half2 a, const __half2 b)
Performs half2 vector less-than comparison.
__host__ __device__ __half2 __hltu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-than comparison.
__host__ __device__ unsigned int __hltu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered less-than comparison.
__host__ __device__ __half2 __hmax2(const __half2 a, const __half2 b)
Calculates half2 vector maximum of two inputs.
__host__ __device__ __half2 __hmax2_nan(const __half2 a, const __half2 b)
Calculates half2 vector maximum of two inputs, NaNs pass through.
__host__ __device__ __half2 __hmin2(const __half2 a, const __half2 b)
Calculates half2 vector minimum of two inputs.
__host__ __device__ __half2 __hmin2_nan(const __half2 a, const __half2 b)
Calculates half2 vector minimum of two inputs, NaNs pass through.
114
Chapter 4. Half Precision Intrinsics
__host__ __device__ __half2 __hne2(const __half2 a, const __half2 b)
Performs half2 vector not-equal comparison.
__host__ __device__ unsigned int __hne2_mask(const __half2 a, const __half2 b)
Performs half2 vector not-equal comparison.
__host__ __device__ __half2 __hneu2(const __half2 a, const __half2 b)
Performs half2 vector unordered not-equal comparison.
__host__ __device__ unsigned int __hneu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered not-equal comparison.
__host__ __device__ bool operator!=(const __half2 &lh, const __half2 &rh)
Performs packed half unordered compare not-equal operation.
__host__ __device__ bool operator<(const __half2 &lh, const __half2 &rh)
Performs packed half ordered less-than compare operation.
__host__ __device__ bool operator<=(const __half2 &lh, const __half2 &rh)
Performs packed half ordered less-or-equal compare operation.
__host__ __device__ bool operator==(const __half2 &lh, const __half2 &rh)
Performs packed half ordered compare equal operation.
__host__ __device__ bool operator>(const __half2 &lh, const __half2 &rh)
Performs packed half ordered greater-than compare operation.
__host__ __device__ bool operator>=(const __half2 &lh, const __half2 &rh)
Performs packed half ordered greater-or-equal compare operation.
4.7.1. Functions
__host__ __device__ bool __hbeq2(const __half2 a, const __half2 b)
Performs half2 vector if-equal comparison and returns boolean true if both half results are
true, boolean false otherwise.
Performs half2 vector if-equal comparison of inputs a and b. The bool result is set to true only
if both half if-equal comparisons evaluate to true, or false otherwise. NaN inputs generate false
results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of if-equal comparison of vectors a and b are true;
▶ false otherwise.
__host__ __device__ bool __hbequ2(const __half2 a, const __half2 b)
Performs half2 vector unordered if-equal comparison and returns boolean true if both half
results are true, boolean false otherwise.
4.7. Half2 Comparison Functions
115
Performs half2 vector if-equal comparison of inputs a and b. The bool result is set to true only
if both half if-equal comparisons evaluate to true, or false otherwise. NaN inputs generate true
results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of unordered if-equal comparison of vectors a and b
are true;
▶ false otherwise.
__host__ __device__ bool __hbge2(const __half2 a, const __half2 b)
Performs half2 vector greater-equal comparison and returns boolean true if both half results
are true, boolean false otherwise.
Performs half2 vector greater-equal comparison of inputs a and b. The bool result is set to true
only if both half greater-equal comparisons evaluate to true, or false otherwise. NaN inputs
generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of greater-equal comparison of vectors a and b are
true;
▶ false otherwise.
__host__ __device__ bool __hbgeu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-equal comparison and returns boolean true if both
half results are true, boolean false otherwise.
Performs half2 vector greater-equal comparison of inputs a and b. The bool result is set to true
only if both half greater-equal comparisons evaluate to true, or false otherwise. NaN inputs
generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of unordered greater-equal comparison of vectors a
and b are true;
▶ false otherwise.
116
Chapter 4. Half Precision Intrinsics
__host__ __device__ bool __hbgt2(const __half2 a, const __half2 b)
Performs half2 vector greater-than comparison and returns boolean true if both half results
are true, boolean false otherwise.
Performs half2 vector greater-than comparison of inputs a and b. The bool result is set to
true only if both half greater-than comparisons evaluate to true, or false otherwise. NaN inputs
generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of greater-than comparison of vectors a and b are true;
▶ false otherwise.
__host__ __device__ bool __hbgtu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-than comparison and returns boolean true if both
half results are true, boolean false otherwise.
Performs half2 vector greater-than comparison of inputs a and b. The bool result is set to
true only if both half greater-than comparisons evaluate to true, or false otherwise. NaN inputs
generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of unordered greater-than comparison of vectors a
and b are true;
▶ false otherwise.
__host__ __device__ bool __hble2(const __half2 a, const __half2 b)
Performs half2 vector less-equal comparison and returns boolean true if both half results are
true, boolean false otherwise.
Performs half2 vector less-equal comparison of inputs a and b. The bool result is set to true only
if both half less-equal comparisons evaluate to true, or false otherwise. NaN inputs generate
false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of less-equal comparison of vectors a and b are true;
4.7. Half2 Comparison Functions
117
▶ false otherwise.
__host__ __device__ bool __hbleu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-equal comparison and returns boolean true if both half
results are true, boolean false otherwise.
Performs half2 vector less-equal comparison of inputs a and b. The bool result is set to true only
if both half less-equal comparisons evaluate to true, or false otherwise. NaN inputs generate
true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of unordered less-equal comparison of vectors a and
b are true;
▶ false otherwise.
__host__ __device__ bool __hblt2(const __half2 a, const __half2 b)
Performs half2 vector less-than comparison and returns boolean true if both half results are
true, boolean false otherwise.
Performs half2 vector less-than comparison of inputs a and b. The bool result is set to true only
if both half less-than comparisons evaluate to true, or false otherwise. NaN inputs generate
false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of less-than comparison of vectors a and b are true;
▶ false otherwise.
__host__ __device__ bool __hbltu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-than comparison and returns boolean true if both half
results are true, boolean false otherwise.
Performs half2 vector less-than comparison of inputs a and b. The bool result is set to true only
if both half less-than comparisons evaluate to true, or false otherwise. NaN inputs generate true
results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
118
Chapter 4. Half Precision Intrinsics
▶ true if both half results of unordered less-than comparison of vectors a and b
are true;
▶ false otherwise.
__host__ __device__ bool __hbne2(const __half2 a, const __half2 b)
Performs half2 vector not-equal comparison and returns boolean true if both half results are
true, boolean false otherwise.
Performs half2 vector not-equal comparison of inputs a and b. The bool result is set to true only
if both half not-equal comparisons evaluate to true, or false otherwise. NaN inputs generate
false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of not-equal comparison of vectors a and b are true,
▶ false otherwise.
__host__ __device__ bool __hbneu2(const __half2 a, const __half2 b)
Performs half2 vector unordered not-equal comparison and returns boolean true if both half
results are true, boolean false otherwise.
Performs half2 vector not-equal comparison of inputs a and b. The bool result is set to true only
if both half not-equal comparisons evaluate to true, or false otherwise. NaN inputs generate true
results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
bool
▶ true if both half results of unordered not-equal comparison of vectors a and b
are true;
▶ false otherwise.
__host__ __device__ __half2 __heq2(const __half2 a, const __half2 b)
Performs half2 vector if-equal comparison.
Performs half2 vector if-equal comparison of inputs a and b. The corresponding half results
are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
4.7. Half2 Comparison Functions
119
▶ The vector result of if-equal comparison of vectors a and b.
__host__ __device__ unsigned int __heq2_mask(const __half2 a, const __half2 b)
Performs half2 vector if-equal comparison.
Performs half2 vector if-equal comparison of inputs a and b. The corresponding unsigned bits
are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of if-equal comparison of vectors a and b.
__host__ __device__ __half2 __hequ2(const __half2 a, const __half2 b)
Performs half2 vector unordered if-equal comparison.
Performs half2 vector if-equal comparison of inputs a and b. The corresponding half results
are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The vector result of unordered if-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hequ2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered if-equal comparison.
Performs half2 vector if-equal comparison of inputs a and b. The corresponding unsigned bits
are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered if-equal comparison of vectors a and b.
__host__ __device__ __half2 __hge2(const __half2 a, const __half2 b)
Performs half2 vector greater-equal comparison.
Performs half2 vector greater-equal comparison of inputs a and b. The corresponding half
results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
120
Chapter 4. Half Precision Intrinsics
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The vector result of greater-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hge2_mask(const __half2 a, const __half2 b)
Performs half2 vector greater-equal comparison.
Performs half2 vector greater-equal comparison of inputs a and b. The corresponding un-
signed bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of greater-equal comparison of vectors a and b.
__host__ __device__ __half2 __hgeu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-equal comparison.
Performs half2 vector greater-equal comparison of inputs a and b. The corresponding half
results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The half2 vector result of unordered greater-equal comparison of vectors a
and b.
__host__ __device__ unsigned int __hgeu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-equal comparison.
Performs half2 vector greater-equal comparison of inputs a and b. The corresponding un-
signed bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered greater-equal comparison of vectors a and
b.
4.7. Half2 Comparison Functions
121
__host__ __device__ __half2 __hgt2(const __half2 a, const __half2 b)
Performs half2 vector greater-than comparison.
Performs half2 vector greater-than comparison of inputs a and b. The corresponding half
results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The vector result of greater-than comparison of vectors a and b.
__host__ __device__ unsigned int __hgt2_mask(const __half2 a, const __half2 b)
Performs half2 vector greater-than comparison.
Performs half2 vector greater-than comparison of inputs a and b. The corresponding unsigned
bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of greater-than comparison of vectors a and b.
__host__ __device__ __half2 __hgtu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-than comparison.
Performs half2 vector greater-than comparison of inputs a and b. The corresponding half
results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The half2 vector result of unordered greater-than comparison of vectors a and
b.
__host__ __device__ unsigned int __hgtu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-than comparison.
Performs half2 vector greater-than comparison of inputs a and b. The corresponding unsigned
bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
122
Chapter 4. Half Precision Intrinsics
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered greater-than comparison of vectors a and
b.
__host__ __device__ __half2 __hisnan2(const __half2 a)
Determine whether half2 argument is a NaN.
Determine whether each half of input half2 number a is a NaN.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The half2 with the corresponding half results set to 1.0 for NaN, 0.0 otherwise.
__host__ __device__ __half2 __hle2(const __half2 a, const __half2 b)
Performs half2 vector less-equal comparison.
Performs half2 vector less-equal comparison of inputs a and b. The corresponding half results
are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The half2 result of less-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hle2_mask(const __half2 a, const __half2 b)
Performs half2 vector less-equal comparison.
Performs half2 vector less-equal comparison of inputs a and b. The corresponding unsigned
bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of less-equal comparison of vectors a and b.
__host__ __device__ __half2 __hleu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-equal comparison.
Performs half2 vector less-equal comparison of inputs a and b. The corresponding half results
are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
4.7. Half2 Comparison Functions
123
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The vector result of unordered less-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hleu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered less-equal comparison.
Performs half2 vector less-equal comparison of inputs a and b. The corresponding unsigned
bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered less-equal comparison of vectors a and b.
__host__ __device__ __half2 __hlt2(const __half2 a, const __half2 b)
Performs half2 vector less-than comparison.
Performs half2 vector less-than comparison of inputs a and b. The corresponding half results
are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The half2 vector result of less-than comparison of vectors a and b.
__host__ __device__ unsigned int __hlt2_mask(const __half2 a, const __half2 b)
Performs half2 vector less-than comparison.
Performs half2 vector less-than comparison of inputs a and b. The corresponding unsigned
bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of less-than comparison of vectors a and b.
124
Chapter 4. Half Precision Intrinsics
__host__ __device__ __half2 __hltu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-than comparison.
Performs half2 vector less-than comparison of inputs a and b. The corresponding half results
are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The vector result of unordered less-than comparison of vectors a and b.
__host__ __device__ unsigned int __hltu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered less-than comparison.
Performs half2 vector less-than comparison of inputs a and b. The corresponding unsigned
bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered less-than comparison of vectors a and b.
__host__ __device__ __half2 __hmax2(const __half2 a, const __half2 b)
Calculates half2 vector maximum of two inputs.
Calculates half2 vector max(a, b). Elementwise half operation is 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
▶ The result of elementwise maximum of vectors a and b
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
__host__ __device__ __half2 __hmax2_nan(const __half2 a, const __half2 b)
Calculates half2 vector maximum of two inputs, NaNs pass through.
Calculates half2 vector max(a, b). Elementwise half operation is defined as (a > b) ? a : b.
▶ If either of inputs is NaN, then canonical NaN is returned.
4.7. Half2 Comparison Functions
125
▶ If values of both inputs are 0.0, then +0.0 > -0.0
▶ The result of elementwise maximum of vectors a and b, with NaNs pass through
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
__host__ __device__ __half2 __hmin2(const __half2 a, const __half2 b)
Calculates half2 vector minimum of two inputs.
Calculates half2 vector min(a, b). Elementwise half operation is 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
▶ The result of elementwise minimum of vectors a and b
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
__host__ __device__ __half2 __hmin2_nan(const __half2 a, const __half2 b)
Calculates half2 vector minimum of two inputs, NaNs pass through.
Calculates half2 vector min(a, b). Elementwise half operation is 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
▶ The result of elementwise minimum of vectors a and b, with NaNs pass through
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
__host__ __device__ __half2 __hne2(const __half2 a, const __half2 b)
Performs half2 vector not-equal comparison.
Performs half2 vector not-equal comparison of inputs a and b. The corresponding half results
are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
126
Chapter 4. Half Precision Intrinsics
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The vector result of not-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hne2_mask(const __half2 a, const __half2 b)
Performs half2 vector not-equal comparison.
Performs half2 vector not-equal comparison of inputs a and b. The corresponding unsigned
bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of not-equal comparison of vectors a and b.
__host__ __device__ __half2 __hneu2(const __half2 a, const __half2 b)
Performs half2 vector unordered not-equal comparison.
Performs half2 vector not-equal comparison of inputs a and b. The corresponding half results
are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
half2
▶ The vector result of unordered not-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hneu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered not-equal comparison.
Performs half2 vector not-equal comparison of inputs a and b. The corresponding unsigned
bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - half2. Is only being read.
▶ b - [in] - half2. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered not-equal comparison of vectors a and b.
4.7. Half2 Comparison Functions
127
__host__ __device__ bool operator!=(const __half2 &lh, const __half2 &rh)
Performs packed half unordered compare not-equal operation.
See also:
__hbneu2(__half2, __half2)
__host__ __device__ bool operator<(const __half2 &lh, const __half2 &rh)
Performs packed half ordered less-than compare operation.
See also:
__hblt2(__half2, __half2)
__host__ __device__ bool operator<=(const __half2 &lh, const __half2 &rh)
Performs packed half ordered less-or-equal compare operation.
See also:
__hble2(__half2, __half2)
__host__ __device__ bool operator==(const __half2 &lh, const __half2 &rh)
Performs packed half ordered compare equal operation.
See also:
__hbeq2(__half2, __half2)
__host__ __device__ bool operator>(const __half2 &lh, const __half2 &rh)
Performs packed half ordered greater-than compare operation.
See also:
__hbgt2(__half2, __half2)
__host__ __device__ bool operator>=(const __half2 &lh, const __half2 &rh)
Performs packed half ordered greater-or-equal compare operation.
See also:
__hbge2(__half2, __half2)
128
Chapter 4. Half Precision Intrinsics
4.8. Half2 Math Functions
To use these functions, include the header file cuda_fp16.h in your program.
Functions
__device__ __half2 h2ceil(const __half2 h)
Calculate half2 vector ceiling of the input argument.
__device__ __half2 h2cos(const __half2 a)
Calculates half2 vector cosine in round-to-nearest-even mode.
__device__ __half2 h2exp(const __half2 a)
Calculates half2 vector exponential function in round-to-nearest-even mode.
__device__ __half2 h2exp10(const __half2 a)
Calculates half2 vector decimal exponential function in round-to-nearest-even mode.
__device__ __half2 h2exp2(const __half2 a)
Calculates half2 vector binary exponential function in round-to-nearest-even mode.
__device__ __half2 h2floor(const __half2 h)
Calculate the largest integer less than or equal to h .
__device__ __half2 h2log(const __half2 a)
Calculates half2 vector natural logarithm in round-to-nearest-even mode.
__device__ __half2 h2log10(const __half2 a)
Calculates half2 vector decimal logarithm in round-to-nearest-even mode.
__device__ __half2 h2log2(const __half2 a)
Calculates half2 vector binary logarithm in round-to-nearest-even mode.
__device__ __half2 h2rcp(const __half2 a)
Calculates half2 vector reciprocal in round-to-nearest-even mode.
__device__ __half2 h2rint(const __half2 h)
Round input to nearest integer value in half-precision floating-point number.
__device__ __half2 h2rsqrt(const __half2 a)
Calculates half2 vector reciprocal square root in round-to-nearest-even mode.
__device__ __half2 h2sin(const __half2 a)
Calculates half2 vector sine in round-to-nearest-even mode.
__device__ __half2 h2sqrt(const __half2 a)
Calculates half2 vector square root in round-to-nearest-even mode.
__device__ __half2 h2tanh(const __half2 a)
Calculates half2 vector hyperbolic tangent function in round-to-nearest-even mode.
__device__ __half2 h2tanh_approx(const __half2 a)
Calculates half2 vector approximate hyperbolic tangent function.
__device__ __half2 h2trunc(const __half2 h)
Truncate half2 vector input argument to the integral part.
4.8. Half2 Math Functions
129
4.8.1. Functions
__device__ __half2 h2ceil(const __half2 h)
Calculate half2 vector ceiling of the input argument.
For each component of vector h compute the smallest integer value not less than h.
See also:
hceil(__half) for further details.
Parameters
h - [in] - half2. Is only being read.
Returns
half2
▶ The vector of smallest integers not less than h.
__device__ __half2 h2cos(const __half2 a)
Calculates half2 vector cosine in round-to-nearest-even mode.
Calculates half2 cosine of input vector a in round-to-nearest-even mode.
See also:
hcos(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise cosine on vector a.
__device__ __half2 h2exp(const __half2 a)
Calculates half2 vector exponential function in round-to-nearest-even mode.
Calculates half2 exponential function of input vector a in round-to-nearest-even mode.
See also:
hexp(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise exponential function on vector a.
__device__ __half2 h2exp10(const __half2 a)
Calculates half2 vector decimal exponential function in round-to-nearest-even mode.
Calculates half2 decimal exponential function of input vector a in round-to-nearest-even mode.
See also:
hexp10(__half) for further details.
130
Chapter 4. Half Precision Intrinsics
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise decimal exponential function on vector a.
__device__ __half2 h2exp2(const __half2 a)
Calculates half2 vector binary exponential function in round-to-nearest-even mode.
Calculates half2 binary exponential function of input vector a in round-to-nearest-even mode.
See also:
hexp2(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise binary exponential function on vector a.
__device__ __half2 h2floor(const __half2 h)
Calculate the largest integer less than or equal to h.
For each component of vector h calculate the largest integer value which is less than or equal to
h.
See also:
hfloor(__half) for further details.
Parameters
h - [in] - half2. Is only being read.
Returns
half2
▶ The vector of largest integers which is less than or equal to h.
__device__ __half2 h2log(const __half2 a)
Calculates half2 vector natural logarithm in round-to-nearest-even mode.
Calculates half2 natural logarithm of input vector a in round-to-nearest-even mode.
See also:
hlog(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise natural logarithm on vector a.
4.8. Half2 Math Functions
131
__device__ __half2 h2log10(const __half2 a)
Calculates half2 vector decimal logarithm in round-to-nearest-even mode.
Calculates half2 decimal logarithm of input vector a in round-to-nearest-even mode.
See also:
hlog10(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise decimal logarithm on vector a.
__device__ __half2 h2log2(const __half2 a)
Calculates half2 vector binary logarithm in round-to-nearest-even mode.
Calculates half2 binary logarithm of input vector a in round-to-nearest-even mode.
See also:
hlog2(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise binary logarithm on vector a.
__device__ __half2 h2rcp(const __half2 a)
Calculates half2 vector reciprocal in round-to-nearest-even mode.
Calculates half2 reciprocal of input vector a in round-to-nearest-even mode.
See also:
hrcp(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise reciprocal on vector a.
__device__ __half2 h2rint(const __half2 h)
Round input to nearest integer value in half-precision floating-point number.
Round each component of half2 vector h to the nearest integer value in half-precision floating-
point format, with halfway cases rounded to the nearest even integer value.
See also:
hrint(__half) for further details.
132
Chapter 4. Half Precision Intrinsics
Parameters
h - [in] - half2. Is only being read.
Returns
half2
▶ The vector of rounded integer values.
__device__ __half2 h2rsqrt(const __half2 a)
Calculates half2 vector reciprocal square root in round-to-nearest-even mode.
Calculates half2 reciprocal square root of input vector a in round-to-nearest-even mode.
See also:
hrsqrt(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise reciprocal square root on vector a.
__device__ __half2 h2sin(const __half2 a)
Calculates half2 vector sine in round-to-nearest-even mode.
Calculates half2 sine of input vector a in round-to-nearest-even mode.
See also:
hsin(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise sine on vector a.
__device__ __half2 h2sqrt(const __half2 a)
Calculates half2 vector square root in round-to-nearest-even mode.
Calculates half2 square root of input vector a in round-to-nearest-even mode.
See also:
hsqrt(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise square root on vector a.
4.8. Half2 Math Functions
133
__device__ __half2 h2tanh(const __half2 a)
Calculates half2 vector hyperbolic tangent function in round-to-nearest-even mode.
Calculates half2 hyperbolic tangent function of input vector a in round-to-nearest-even mode.
See also:
htanh(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise hyperbolic tangent function on vector a.
__device__ __half2 h2tanh_approx(const __half2 a)
Calculates half2 vector approximate hyperbolic tangent function.
Calculates half2 approximate hyperbolic tangent function of input vector a. This operation uses
HW acceleration on devices of compute capability 7.5 and higher.
See also:
htanh_approx(__half) for further details.
Parameters
a - [in] - half2. Is only being read.
Returns
half2
▶ The elementwise approximate hyperbolic tangent function on vector a.
__device__ __half2 h2trunc(const __half2 h)
Truncate half2 vector input argument to the integral part.
Round each component of vector h to the largest integer value that does not exceed h in mag-
nitude.
See also:
htrunc(__half) for further details.
Parameters
h - [in] - half2. Is only being read.
Returns
half2
▶ The truncated h.
134
Chapter 4. Half Precision Intrinsics
Groups
Half Arithmetic Constants
To use these constants, include the header file cuda_fp16.h in your program.
Half Arithmetic Functions
To use these functions, include the header file cuda_fp16.h in your program.
Half Comparison Functions
To use these functions, include the header file cuda_fp16.h in your program.
Half Math Functions
To use these functions, include the header file cuda_fp16.h in your program.
Half Precision Conversion and Data Movement
To use these functions, include the header file cuda_fp16.h in your program.
Half2 Arithmetic Functions
To use these functions, include the header file cuda_fp16.h in your program.
Half2 Comparison Functions
To use these functions, include the header file cuda_fp16.h in your program.
Half2 Math Functions
To use these functions, include the header file cuda_fp16.h in your program.
Structs
__half
__half data type
__half2
__half2 data type
__half2_raw
__half2_raw data type
__half_raw
__half_raw data type
Typedefs
__nv_half
This datatype is an __nv_ prefixed alias.
__nv_half2
This datatype is an __nv_ prefixed alias.
__nv_half2_raw
This datatype is an __nv_ prefixed alias.
__nv_half_raw
This datatype is an __nv_ prefixed alias.
half This datatype is meant to be the first-class or fundamental implementation of the half-precision
numbers format.
4.8. Half2 Math Functions
135
half2
This datatype is meant to be the first-class or fundamental implementation of type for pairs of
half-precision numbers.
nv_half
This datatype is an nv_ prefixed alias.
nv_half2
This datatype is an nv_ prefixed alias.
4.9. Typedefs
typedef __half __nv_half
This datatype is an __nv_ prefixed alias.
typedef __half2 __nv_half2
This datatype is an __nv_ prefixed alias.
typedef __half2_raw __nv_half2_raw
This datatype is an __nv_ prefixed alias.
typedef __half_raw __nv_half_raw
This datatype is an __nv_ prefixed alias.
typedef __half half
This datatype is meant to be the first-class or fundamental implementation of the half-precision
numbers format.
Should be implemented in the compiler in the future. Current implementation is a simple typedef
to a respective user-level type with underscores.
typedef __half2 half2
This datatype is meant to be the first-class or fundamental implementation of type for pairs of
half-precision numbers.
Should be implemented in the compiler in the future. Current implementation is a simple typedef
to a respective user-level type with underscores.
typedef __half nv_half
This datatype is an nv_ prefixed alias.
typedef __half2 nv_half2
This datatype is an nv_ prefixed alias.
136
Chapter 4. Half Precision Intrinsics
Chapter 5. Bfloat16 Precision Intrinsics
This section describes nv_bfloat16 precision intrinsic functions.
To use these functions, include the header file cuda_bf16.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. Specific examples are:
▶ hsin(__nv_bfloat16);
▶ hcos(__nv_bfloat16);
▶ h2sin(__nv_bfloat162);
▶ h2cos(__nv_bfloat162);
The following macros are available to help users selectively enable/disable various definitions present
in the header file:
▶ CUDA_NO_BFLOAT16 - 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_BFLOAT16_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 __nv_bfloat16 which is essentially a user-
defined type.
▶ __CUDA_NO_BFLOAT16_OPERATORS__ and __CUDA_NO_BFLOAT162_OPERATORS__ - If defined,
these macros will prevent the inadvertent use of usual arithmetic and comparison opera-
tors. This enforces the storage-only type semantics and prevents C++ style computations on
__nv_bfloat16 and __nv_bfloat162 types.
5.1. Bfloat16 Arithmetic Constants
To use these constants, include the header file cuda_bf16.h in your program.
137
Macros
CUDART_INF_BF16
Defines floating-point positive infinity value for the nv_bfloat16 data type.
CUDART_MAX_NORMAL_BF16
Defines a maximum representable value for the nv_bfloat16 data type.
CUDART_MIN_DENORM_BF16
Defines a minimum representable (denormalized) value for the nv_bfloat16 data type.
CUDART_NAN_BF16
Defines canonical NaN value for the nv_bfloat16 data type.
CUDART_NEG_ZERO_BF16
Defines a negative zero value for the nv_bfloat16 data type.
CUDART_ONE_BF16
Defines a value of 1.0 for the nv_bfloat16 data type.
CUDART_ZERO_BF16
Defines a positive zero value for the nv_bfloat16 data type.
5.1.1. Macros
CUDART_INF_BF16 __ushort_as_bfloat16((unsigned short)0x7F80U)
Defines floating-point positive infinity value for the nv_bfloat16 data type.
CUDART_MAX_NORMAL_BF16 __ushort_as_bfloat16((unsigned short)0x7F7FU)
Defines a maximum representable value for the nv_bfloat16 data type.
CUDART_MIN_DENORM_BF16 __ushort_as_bfloat16((unsigned short)0x0001U)
Defines a minimum representable (denormalized) value for the nv_bfloat16 data type.
CUDART_NAN_BF16 __ushort_as_bfloat16((unsigned short)0x7FFFU)
Defines canonical NaN value for the nv_bfloat16 data type.
CUDART_NEG_ZERO_BF16 __ushort_as_bfloat16((unsigned short)0x8000U)
Defines a negative zero value for the nv_bfloat16 data type.
CUDART_ONE_BF16 __ushort_as_bfloat16((unsigned short)0x3F80U)
Defines a value of 1.0 for the nv_bfloat16 data type.
CUDART_ZERO_BF16 __ushort_as_bfloat16((unsigned short)0x0000U)
Defines a positive zero value for the nv_bfloat16 data type.
138
Chapter 5. Bfloat16 Precision Intrinsics
5.2. Bfloat16 Arithmetic Functions
To use these functions, include the header file cuda_bf16.h in your program.
Functions
__host__ __device__ __nv_bfloat16 __habs(const __nv_bfloat16 a)
Calculates the absolute value of input nv_bfloat16 number and returns the result.
__host__ __device__ __nv_bfloat16 __hadd(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 addition in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat16 __hadd_rn(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 addition in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat16 __hadd_sat(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 addition in round-to-nearest-even mode, with saturation to [0.0, 1.0].
__host__ __device__ __nv_bfloat16 __hdiv(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 division in round-to-nearest-even mode.
__device__ __nv_bfloat16 __hfma(const __nv_bfloat16 a, const __nv_bfloat16 b, const __nv_bfloat16
c)
Performs nv_bfloat16 fused multiply-add in round-to-nearest-even mode.
__device__
__nv_bfloat16
__hfma_relu(const
__nv_bfloat16 a, const
__nv_bfloat16 b, const
__nv_bfloat16 c)
Performs nv_bfloat16 fused multiply-add in round-to-nearest-even mode with relu saturation.
__device__
__nv_bfloat16
__hfma_sat(const
__nv_bfloat16 a, const
__nv_bfloat16 b, const
__nv_bfloat16 c)
Performs nv_bfloat16 fused multiply-add in round-to-nearest-even mode, with saturation to
[0.0, 1.0].
__host__ __device__ __nv_bfloat16 __hmul(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 multiplication in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat16 __hmul_rn(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 multiplication in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat16 __hmul_sat(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 multiplication in round-to-nearest-even mode, with saturation to [0.0,
1.0].
__host__ __device__ __nv_bfloat16 __hneg(const __nv_bfloat16 a)
Negates input nv_bfloat16 number and returns the result.
__host__ __device__ __nv_bfloat16 __hsub(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 subtraction in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat16 __hsub_rn(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 subtraction in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat16 __hsub_sat(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 subtraction in round-to-nearest-even mode, with saturation to [0.0, 1.0].
__device__ __nv_bfloat16 atomicAdd(__nv_bfloat16 *const address, const __nv_bfloat16 val)
Adds val to the value stored at address in global or shared memory, and writes this value back
to address .
5.2. Bfloat16 Arithmetic Functions
139
__host__ __device__ __nv_bfloat16 operator*(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 multiplication operation.
__host__ __device__ __nv_bfloat16 & operator*=(__nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 compound assignment with multiplication operation.
__host__ __device__ __nv_bfloat16 operator+(const __nv_bfloat16 &h)
Implements nv_bfloat16 unary plus operator, returns input value.
__host__ __device__ __nv_bfloat16 operator+(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 addition operation.
__host__ __device__ __nv_bfloat16 operator++(__nv_bfloat16 &h, const int ignored)
Performs nv_bfloat16 postfix increment operation.
__host__ __device__ __nv_bfloat16 & operator++(__nv_bfloat16 &h)
Performs nv_bfloat16 prefix increment operation.
__host__ __device__ __nv_bfloat16 & operator+=(__nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 compound assignment with addition operation.
__host__ __device__ __nv_bfloat16 operator-(const __nv_bfloat16 &h)
Implements nv_bfloat16 unary minus operator.
__host__ __device__ __nv_bfloat16 operator-(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 subtraction operation.
__host__ __device__ __nv_bfloat16 & operator-(__nv_bfloat16 &h)
Performs nv_bfloat16 prefix decrement operation.
__host__ __device__ __nv_bfloat16 operator-(__nv_bfloat16 &h, const int ignored)
Performs nv_bfloat16 postfix decrement operation.
__host__ __device__ __nv_bfloat16 & operator-=(__nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 compound assignment with subtraction operation.
__host__ __device__ __nv_bfloat16 operator/(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 division operation.
__host__ __device__ __nv_bfloat16 & operator/=(__nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 compound assignment with division operation.
5.2.1. Functions
__host__ __device__ __nv_bfloat16 __habs(const __nv_bfloat16 a)
Calculates the absolute value of input nv_bfloat16 number and returns the result.
Calculates the absolute value of input nv_bfloat16 number and returns the result.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The absolute value of a.
__host__ __device__ __nv_bfloat16 __hadd(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 addition in round-to-nearest-even mode.
Performs nv_bfloat16 addition of inputs a and b, in round-to-nearest-even mode.
140
Chapter 5. Bfloat16 Precision Intrinsics
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The sum of a and b.
__host__ __device__ __nv_bfloat16 __hadd_rn(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 addition in round-to-nearest-even mode.
Performs nv_bfloat16 addition of inputs a and b, in round-to-nearest-even mode. Prevents
floating-point contractions of mul+add into fma.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The sum of a and b.
__host__ __device__ __nv_bfloat16 __hadd_sat(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 addition in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs nv_bfloat16 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] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The sum of a and b, with respect to saturation.
__host__ __device__ __nv_bfloat16 __hdiv(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 division in round-to-nearest-even mode.
Divides nv_bfloat16 input a by input b in round-to-nearest-even mode.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The result of dividing a by b.
5.2. Bfloat16 Arithmetic Functions
141
__device__ __nv_bfloat16 __hfma(const __nv_bfloat16 a, const __nv_bfloat16 b, const __nv_bfloat16
c)
Performs nv_bfloat16 fused multiply-add in round-to-nearest-even mode.
Performs nv_bfloat16 multiply on inputs a and b, then performs a nv_bfloat16 add of the
result with c, rounding the result once in round-to-nearest-even mode.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
▶ c - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The result of fused multiply-add operation on a, b, and c.
__device__ __nv_bfloat16 __hfma_relu(const __nv_bfloat16 a, const __nv_bfloat16 b, const
__nv_bfloat16 c)
Performs nv_bfloat16 fused multiply-add in round-to-nearest-even mode with relu saturation.
Performs nv_bfloat16 multiply on inputs a and b, then performs a nv_bfloat16 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] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
▶ c - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The result of fused multiply-add operation on a, b, and c with relu saturation.
__device__ __nv_bfloat16 __hfma_sat(const __nv_bfloat16 a, const __nv_bfloat16 b, const
__nv_bfloat16 c)
Performs nv_bfloat16 fused multiply-add in round-to-nearest-even mode, with saturation to
[0.0, 1.0].
Performs nv_bfloat16 multiply on inputs a and b, then performs a nv_bfloat16 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] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
▶ c - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
142
Chapter 5. Bfloat16 Precision Intrinsics
▶ The result of fused multiply-add operation on a, b, and c, with respect to satu-
ration.
__host__ __device__ __nv_bfloat16 __hmul(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 multiplication in round-to-nearest-even mode.
Performs nv_bfloat16 multiplication of inputs a and b, in round-to-nearest-even mode.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The result of multiplying a and b.
__host__ __device__ __nv_bfloat16 __hmul_rn(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 multiplication in round-to-nearest-even mode.
Performs nv_bfloat16 multiplication of inputs a and b, in round-to-nearest-even mode. Pre-
vents floating-point contractions of mul+add or sub into fma.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The result of multiplying a and b.
__host__ __device__ __nv_bfloat16 __hmul_sat(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 multiplication in round-to-nearest-even mode, with saturation to [0.0,
1.0].
Performs nv_bfloat16 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] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The result of multiplying a and b, with respect to saturation.
__host__ __device__ __nv_bfloat16 __hneg(const __nv_bfloat16 a)
Negates input nv_bfloat16 number and returns the result.
Negates input nv_bfloat16 number and returns the result.
Parameters
a - [in] - nv_bfloat16. Is only being read.
5.2. Bfloat16 Arithmetic Functions
143
Returns
nv_bfloat16
▶ minus a
__host__ __device__ __nv_bfloat16 __hsub(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 subtraction in round-to-nearest-even mode.
Subtracts nv_bfloat16 input b from input a in round-to-nearest-even mode.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The result of subtracting b from a.
__host__ __device__ __nv_bfloat16 __hsub_rn(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 subtraction in round-to-nearest-even mode.
Subtracts nv_bfloat16 input b from input a in round-to-nearest-even mode. Prevents floating-
point contractions of mul+sub into fma.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The result of subtracting b from a.
__host__ __device__ __nv_bfloat16 __hsub_sat(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 subtraction in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Subtracts nv_bfloat16 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] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The result of subtraction of b from a, with respect to saturation.
__device__ __nv_bfloat16 atomicAdd(__nv_bfloat16 *const address, const __nv_bfloat16 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.
144
Chapter 5. Bfloat16 Precision Intrinsics
The location of address must be in global or shared memory. This operation has undefined
behavior otherwise. This operation is natively supported by devices of compute capability 9.x
and higher, older devices of compute capability 7.x and 8.x use emulation path.
Note: For more details about this function, see the Atomic Functions section in the CUDA C++
Programming Guide.
Parameters
▶ address - [in] - __nv_bfloat16*. An address in global or shared memory.
▶ val - [in] - __nv_bfloat16. The value to be added.
Returns
__nv_bfloat16
▶ The old value read from address.
__host__ __device__ __nv_bfloat16 operator*(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 multiplication operation.
See also __hmul(__nv_bfloat16, __nv_bfloat16)
__host__ __device__ __nv_bfloat16 &operator*=(__nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 compound assignment with multiplication operation.
__host__ __device__ __nv_bfloat16 operator+(const __nv_bfloat16 &h)
Implements nv_bfloat16 unary plus operator, returns input value.
__host__ __device__ __nv_bfloat16 operator+(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 addition operation.
See also __hadd(__nv_bfloat16, __nv_bfloat16)
__host__ __device__ __nv_bfloat16 operator++(__nv_bfloat16 &h, const int ignored)
Performs nv_bfloat16 postfix increment operation.
__host__ __device__ __nv_bfloat16 &operator++(__nv_bfloat16 &h)
Performs nv_bfloat16 prefix increment operation.
__host__ __device__ __nv_bfloat16 &operator+=(__nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 compound assignment with addition operation.
__host__ __device__ __nv_bfloat16 operator-(const __nv_bfloat16 &h)
Implements nv_bfloat16 unary minus operator.
See also __hneg(__nv_bfloat16)
__host__ __device__ __nv_bfloat16 operator-(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 subtraction operation.
See also __hsub(__nv_bfloat16, __nv_bfloat16)
__host__ __device__ __nv_bfloat16 &operator--(__nv_bfloat16 &h)
Performs nv_bfloat16 prefix decrement operation.
5.2. Bfloat16 Arithmetic Functions
145
__host__ __device__ __nv_bfloat16 operator--(__nv_bfloat16 &h, const int ignored)
Performs nv_bfloat16 postfix decrement operation.
__host__ __device__ __nv_bfloat16 &operator-=(__nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 compound assignment with subtraction operation.
__host__ __device__ __nv_bfloat16 operator∕(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 division operation.
See also __hdiv(__nv_bfloat16, __nv_bfloat16)
__host__ __device__ __nv_bfloat16 &operator∕=(__nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 compound assignment with division operation.
5.3. Bfloat16 Comparison Functions
To use these functions, include the header file cuda_bf16.h in your program.
Functions
__host__ __device__ bool __heq(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 if-equal comparison.
__host__ __device__ bool __hequ(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered if-equal comparison.
__host__ __device__ bool __hge(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 greater-equal comparison.
__host__ __device__ bool __hgeu(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered greater-equal comparison.
__host__ __device__ bool __hgt(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 greater-than comparison.
__host__ __device__ bool __hgtu(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered greater-than comparison.
__host__ __device__ int __hisinf(const __nv_bfloat16 a)
Checks if the input nv_bfloat16 number is infinite.
__host__ __device__ bool __hisnan(const __nv_bfloat16 a)
Determine whether nv_bfloat16 argument is a NaN.
__host__ __device__ bool __hle(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 less-equal comparison.
__host__ __device__ bool __hleu(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered less-equal comparison.
__host__ __device__ bool __hlt(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 less-than comparison.
__host__ __device__ bool __hltu(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered less-than comparison.
146
Chapter 5. Bfloat16 Precision Intrinsics
__host__ __device__ __nv_bfloat16 __hmax(const __nv_bfloat16 a, const __nv_bfloat16 b)
Calculates nv_bfloat16 maximum of two input values.
__host__ __device__ __nv_bfloat16 __hmax_nan(const __nv_bfloat16 a, const __nv_bfloat16 b)
Calculates nv_bfloat16 maximum of two input values, NaNs pass through.
__host__ __device__ __nv_bfloat16 __hmin(const __nv_bfloat16 a, const __nv_bfloat16 b)
Calculates nv_bfloat16 minimum of two input values.
__host__ __device__ __nv_bfloat16 __hmin_nan(const __nv_bfloat16 a, const __nv_bfloat16 b)
Calculates nv_bfloat16 minimum of two input values, NaNs pass through.
__host__ __device__ bool __hne(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 not-equal comparison.
__host__ __device__ bool __hneu(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered not-equal comparison.
__host__ __device__ bool operator!=(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 unordered compare not-equal operation.
__host__ __device__ bool operator<(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 ordered less-than compare operation.
__host__ __device__ bool operator<=(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 ordered less-or-equal compare operation.
__host__ __device__ bool operator==(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 ordered compare equal operation.
__host__ __device__ bool operator>(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 ordered greater-than compare operation.
__host__ __device__ bool operator>=(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 ordered greater-or-equal compare operation.
5.3.1. Functions
__host__ __device__ bool __heq(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 if-equal comparison.
Performs nv_bfloat16 if-equal comparison of inputs a and b. NaN inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of if-equal comparison of a and b.
__host__ __device__ bool __hequ(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered if-equal comparison.
Performs nv_bfloat16 if-equal comparison of inputs a and b. NaN inputs generate true results.
5.3. Bfloat16 Comparison Functions
147
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of unordered if-equal comparison of a and b.
__host__ __device__ bool __hge(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 greater-equal comparison.
Performs nv_bfloat16 greater-equal comparison of inputs a and b. NaN inputs generate false
results.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of greater-equal comparison of a and b.
__host__ __device__ bool __hgeu(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered greater-equal comparison.
Performs nv_bfloat16 greater-equal comparison of inputs a and b. NaN inputs generate true
results.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of unordered greater-equal comparison of a and b.
__host__ __device__ bool __hgt(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 greater-than comparison.
Performs nv_bfloat16 greater-than comparison of inputs a and b. NaN inputs generate false
results.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of greater-than comparison of a and b.
148
Chapter 5. Bfloat16 Precision Intrinsics
__host__ __device__ bool __hgtu(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered greater-than comparison.
Performs nv_bfloat16 greater-than comparison of inputs a and b. NaN inputs generate true
results.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of unordered greater-than comparison of a and b.
__host__ __device__ int __hisinf(const __nv_bfloat16 a)
Checks if the input nv_bfloat16 number is infinite.
Checks if the input nv_bfloat16 number a is infinite.
Parameters
a - [in] - nv_bfloat16. Is only being read.
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 __nv_bfloat16 a)
Determine whether nv_bfloat16 argument is a NaN.
Determine whether nv_bfloat16 value a is a NaN.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ true if argument is NaN.
__host__ __device__ bool __hle(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 less-equal comparison.
Performs nv_bfloat16 less-equal comparison of inputs a and b. NaN inputs generate false
results.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of less-equal comparison of a and b.
5.3. Bfloat16 Comparison Functions
149
__host__ __device__ bool __hleu(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered less-equal comparison.
Performs nv_bfloat16 less-equal comparison of inputs a and b. NaN inputs generate true re-
sults.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of unordered less-equal comparison of a and b.
__host__ __device__ bool __hlt(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 less-than comparison.
Performs nv_bfloat16 less-than comparison of inputs a and b. NaN inputs generate false re-
sults.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of less-than comparison of a and b.
__host__ __device__ bool __hltu(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered less-than comparison.
Performs nv_bfloat16 less-than comparison of inputs a and b. NaN inputs generate true re-
sults.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of unordered less-than comparison of a and b.
__host__ __device__ __nv_bfloat16 __hmax(const __nv_bfloat16 a, const __nv_bfloat16 b)
Calculates nv_bfloat16 maximum of two input values.
Calculates nv_bfloat16 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
150
Chapter 5. Bfloat16 Precision Intrinsics
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
__host__ __device__ __nv_bfloat16 __hmax_nan(const __nv_bfloat16 a, const __nv_bfloat16 b)
Calculates nv_bfloat16 maximum of two input values, NaNs pass through.
Calculates nv_bfloat16 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] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
__host__ __device__ __nv_bfloat16 __hmin(const __nv_bfloat16 a, const __nv_bfloat16 b)
Calculates nv_bfloat16 minimum of two input values.
Calculates nv_bfloat16 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] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
__host__ __device__ __nv_bfloat16 __hmin_nan(const __nv_bfloat16 a, const __nv_bfloat16 b)
Calculates nv_bfloat16 minimum of two input values, NaNs pass through.
Calculates nv_bfloat16 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] - nv_bfloat16. Is only being read.
5.3. Bfloat16 Comparison Functions
151
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
__host__ __device__ bool __hne(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 not-equal comparison.
Performs nv_bfloat16 not-equal comparison of inputs a and b. NaN inputs generate false re-
sults.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of not-equal comparison of a and b.
__host__ __device__ bool __hneu(const __nv_bfloat16 a, const __nv_bfloat16 b)
Performs nv_bfloat16 unordered not-equal comparison.
Performs nv_bfloat16 not-equal comparison of inputs a and b. NaN inputs generate true re-
sults.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
bool
▶ The boolean result of unordered not-equal comparison of a and b.
__host__ __device__ bool operator!=(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 unordered compare not-equal operation.
See also __hneu(__nv_bfloat16, __nv_bfloat16)
__host__ __device__ bool operator<(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 ordered less-than compare operation.
See also __hlt(__nv_bfloat16, __nv_bfloat16)
__host__ __device__ bool operator<=(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 ordered less-or-equal compare operation.
See also __hle(__nv_bfloat16, __nv_bfloat16)
__host__ __device__ bool operator==(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 ordered compare equal operation.
See also __heq(__nv_bfloat16, __nv_bfloat16)
152
Chapter 5. Bfloat16 Precision Intrinsics
__host__ __device__ bool operator>(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 ordered greater-than compare operation.
See also __hgt(__nv_bfloat16, __nv_bfloat16)
__host__ __device__ bool operator>=(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh)
Performs nv_bfloat16 ordered greater-or-equal compare operation.
See also __hge(__nv_bfloat16, __nv_bfloat16)
5.4. Bfloat16 Math Functions
To use these functions, include the header file cuda_bf16.h in your program.
Functions
__device__ __nv_bfloat16 hceil(const __nv_bfloat16 h)
Calculate ceiling of the input argument.
__device__ __nv_bfloat16 hcos(const __nv_bfloat16 a)
Calculates nv_bfloat16 cosine in round-to-nearest-even mode.
__device__ __nv_bfloat16 hexp(const __nv_bfloat16 a)
Calculates nv_bfloat16 natural exponential function in round-to-nearest-even mode.
__device__ __nv_bfloat16 hexp10(const __nv_bfloat16 a)
Calculates nv_bfloat16 decimal exponential function in round-to-nearest-even mode.
__device__ __nv_bfloat16 hexp2(const __nv_bfloat16 a)
Calculates nv_bfloat16 binary exponential function in round-to-nearest-even mode.
__device__ __nv_bfloat16 hfloor(const __nv_bfloat16 h)
Calculate the largest integer less than or equal to h .
__device__ __nv_bfloat16 hlog(const __nv_bfloat16 a)
Calculates nv_bfloat16 natural logarithm in round-to-nearest-even mode.
__device__ __nv_bfloat16 hlog10(const __nv_bfloat16 a)
Calculates nv_bfloat16 decimal logarithm in round-to-nearest-even mode.
__device__ __nv_bfloat16 hlog2(const __nv_bfloat16 a)
Calculates nv_bfloat16 binary logarithm in round-to-nearest-even mode.
__device__ __nv_bfloat16 hrcp(const __nv_bfloat16 a)
Calculates nv_bfloat16 reciprocal in round-to-nearest-even mode.
__device__ __nv_bfloat16 hrint(const __nv_bfloat16 h)
Round input to nearest integer value in nv_bfloat16 floating-point number.
__device__ __nv_bfloat16 hrsqrt(const __nv_bfloat16 a)
Calculates nv_bfloat16 reciprocal square root in round-to-nearest-even mode.
__device__ __nv_bfloat16 hsin(const __nv_bfloat16 a)
Calculates nv_bfloat16 sine in round-to-nearest-even mode.
__device__ __nv_bfloat16 hsqrt(const __nv_bfloat16 a)
Calculates nv_bfloat16 square root in round-to-nearest-even mode.
5.4. Bfloat16 Math Functions
153
__device__ __nv_bfloat16 htanh(const __nv_bfloat16 a)
Calculates nv_bfloat16 hyperbolic tangent function in round-to-nearest-even mode.
__device__ __nv_bfloat16 htanh_approx(const __nv_bfloat16 a)
Calculates approximate nv_bfloat16 hyperbolic tangent function.
__device__ __nv_bfloat16 htrunc(const __nv_bfloat16 h)
Truncate input argument to the integral part.
5.4.1. Functions
__device__ __nv_bfloat16 hceil(const __nv_bfloat16 h)
Calculate ceiling of the input argument.
Compute the smallest integer value not less than h.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The smallest integer value not less than h.
__device__ __nv_bfloat16 hcos(const __nv_bfloat16 a)
Calculates nv_bfloat16 cosine in round-to-nearest-even mode.
Calculates nv_bfloat16 cosine of input a in round-to-nearest-even mode.
NOTE: this function’s implementation calls cosf(float) function and is exposed to compiler opti-
mizations. Specifically, --use_fast_math flag changes cosf(float) into an intrinsic __cosf(float),
which has less accurate numeric behavior.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The cosine of a.
__device__ __nv_bfloat16 hexp(const __nv_bfloat16 a)
Calculates nv_bfloat16 natural exponential function in round-to-nearest-even mode.
Calculates nv_bfloat16 natural exponential function of input a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The natural exponential function on a.
__device__ __nv_bfloat16 hexp10(const __nv_bfloat16 a)
Calculates nv_bfloat16 decimal exponential function in round-to-nearest-even mode.
Calculates nv_bfloat16 decimal exponential function of input a in round-to-nearest-even
mode.
154
Chapter 5. Bfloat16 Precision Intrinsics
|
||
|
|
|