|
|
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The decimal exponential function on a.
__device__ __nv_bfloat16 hexp2(const __nv_bfloat16 a)
Calculates nv_bfloat16 binary exponential function in round-to-nearest-even mode.
Calculates nv_bfloat16 binary exponential function of input a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The binary exponential function on a.
__device__ __nv_bfloat16 hfloor(const __nv_bfloat16 h)
Calculate the largest integer less than or equal to h.
Calculate the largest integer value which is less than or equal to h.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The largest integer value which is 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.
Calculates nv_bfloat16 natural logarithm of input a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The natural logarithm of a.
__device__ __nv_bfloat16 hlog10(const __nv_bfloat16 a)
Calculates nv_bfloat16 decimal logarithm in round-to-nearest-even mode.
Calculates nv_bfloat16 decimal logarithm of input a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The decimal logarithm of a.
__device__ __nv_bfloat16 hlog2(const __nv_bfloat16 a)
Calculates nv_bfloat16 binary logarithm in round-to-nearest-even mode.
Calculates nv_bfloat16 binary logarithm of input a in round-to-nearest-even mode.
5.4. Bfloat16 Math Functions
155
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The binary logarithm of a.
__device__ __nv_bfloat16 hrcp(const __nv_bfloat16 a)
Calculates nv_bfloat16 reciprocal in round-to-nearest-even mode.
Calculates nv_bfloat16 reciprocal of input a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The reciprocal of a.
__device__ __nv_bfloat16 hrint(const __nv_bfloat16 h)
Round input to nearest integer value in nv_bfloat16 floating-point number.
Round h to the nearest integer value in nv_bfloat16 floating-point format, with bfloat16way cases
rounded to the nearest even integer value.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The nearest integer to h.
__device__ __nv_bfloat16 hrsqrt(const __nv_bfloat16 a)
Calculates nv_bfloat16 reciprocal square root in round-to-nearest-even mode.
Calculates nv_bfloat16 reciprocal square root of input a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The reciprocal square root of a.
__device__ __nv_bfloat16 hsin(const __nv_bfloat16 a)
Calculates nv_bfloat16 sine in round-to-nearest-even mode.
Calculates nv_bfloat16 sine of input a in round-to-nearest-even mode.
NOTE: this function’s implementation calls sinf(float) function and is exposed to compiler opti-
mizations. Specifically, --use_fast_math flag changes sinf(float) into an intrinsic __sinf(float),
which has less accurate numeric behavior.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The sine of a.
156
Chapter 5. Bfloat16 Precision Intrinsics
__device__ __nv_bfloat16 hsqrt(const __nv_bfloat16 a)
Calculates nv_bfloat16 square root in round-to-nearest-even mode.
Calculates nv_bfloat16 square root of input a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The square root of a.
__device__ __nv_bfloat16 htanh(const __nv_bfloat16 a)
Calculates nv_bfloat16 hyperbolic tangent function in round-to-nearest-even mode.
Calculates nv_bfloat16 hyperbolic tangent function: tanh(a) in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The hyperbolic tangent function of a.
▶ htanh (±0) returns (±0).
▶ htanh (±∞) returns (±1).
▶ htanh(NaN) returns NaN.
__device__ __nv_bfloat16 htanh_approx(const __nv_bfloat16 a)
Calculates approximate nv_bfloat16 hyperbolic tangent function.
Calculates approximate nv_bfloat16 hyperbolic tangent function: tanh(a). This operation uses
HW acceleration on devices of compute capability 9.x and higher.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The approximate hyperbolic tangent function of a.
▶ htanh_approx (±0) returns (±0).
▶ htanh_approx (±∞) returns (±1).
▶ htanh_approx(NaN) returns NaN.
__device__ __nv_bfloat16 htrunc(const __nv_bfloat16 h)
Truncate input argument to the integral part.
Round h to the nearest integer value that does not exceed h in magnitude.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat16
▶ The truncated integer value.
5.4. Bfloat16 Math Functions
157
5.5. Bfloat16 Precision Conversion and Data
Movement
To use these functions, include the header file cuda_bf16.h in your program.
Functions
__host__ __device__ float2 __bfloat1622float2(const __nv_bfloat162 a)
Converts both halves of nv_bfloat162 to float2 and returns the result.
__host__ __device__ __nv_bfloat162 __bfloat162bfloat162(const __nv_bfloat16 a)
Returns nv_bfloat162 with both halves equal to the input value.
__host__ __device__ signed char __bfloat162char_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed char in round-towards-zero mode.
__host__ __device__ float __bfloat162float(const __nv_bfloat16 a)
Converts nv_bfloat16 number to float.
__device__ int __bfloat162int_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed integer in round-down mode.
__device__ int __bfloat162int_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed integer in round-to-nearest-even mode.
__device__ int __bfloat162int_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed integer in round-up mode.
__host__ __device__ int __bfloat162int_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed integer in round-towards-zero mode.
__device__ long long int __bfloat162ll_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed 64-bit integer in round-down mode.
__device__ long long int __bfloat162ll_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed 64-bit integer in round-to-nearest-even mode.
__device__ long long int __bfloat162ll_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed 64-bit integer in round-up mode.
__host__ __device__ long long int __bfloat162ll_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed 64-bit integer in round-towards-zero mode.
__device__ short int __bfloat162short_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed short integer in round-down mode.
__device__ short int __bfloat162short_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed short integer in round-to-nearest-even mode.
__device__ short int __bfloat162short_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed short integer in round-up mode.
__host__ __device__ short int __bfloat162short_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed short integer in round-towards-zero mode.
__host__ __device__ unsigned char __bfloat162uchar_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned char in round-towards-zero mode.
158
Chapter 5. Bfloat16 Precision Intrinsics
__device__ unsigned int __bfloat162uint_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned integer in round-down mode.
__device__ unsigned int __bfloat162uint_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned integer in round-to-nearest-even mode.
__device__ unsigned int __bfloat162uint_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned integer in round-up mode.
__host__ __device__ unsigned int __bfloat162uint_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned integer in round-towards-zero mode.
__device__ unsigned long long int __bfloat162ull_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned 64-bit integer in round-down mode.
__device__ unsigned long long int __bfloat162ull_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned 64-bit integer in round-to-nearest-even mode.
__device__ unsigned long long int __bfloat162ull_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned 64-bit integer in round-up mode.
__host__ __device__ unsigned long long int __bfloat162ull_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned 64-bit integer in round-towards-zero mode.
__device__ unsigned short int __bfloat162ushort_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned short integer in round-down mode.
__device__ unsigned short int __bfloat162ushort_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned short integer in round-to-nearest-even mode.
__device__ unsigned short int __bfloat162ushort_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned short integer in round-up mode.
__host__ __device__ unsigned short int __bfloat162ushort_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned short integer in round-towards-zero mode.
__host__ __device__ short int __bfloat16_as_short(const __nv_bfloat16 h)
Reinterprets bits in a nv_bfloat16 as a signed short integer.
__host__ __device__ unsigned short int __bfloat16_as_ushort(const __nv_bfloat16 h)
Reinterprets bits in a nv_bfloat16 as an unsigned short integer.
__host__ __device__ __nv_bfloat16 __double2bfloat16(const double a)
Converts double number to nv_bfloat16 precision in round-to-nearest-even mode and returns
nv_bfloat16 with converted value.
__host__ __device__ __nv_bfloat162 __float22bfloat162_rn(const float2 a)
Converts both components of float2 number to nv_bfloat16 precision in round-to-nearest-even
mode and returns nv_bfloat162 with converted values.
__host__ __device__ __nv_bfloat16 __float2bfloat16(const float a)
Converts float number to nv_bfloat16 precision in round-to-nearest-even mode and returns
nv_bfloat16 with converted value.
__host__ __device__ __nv_bfloat162 __float2bfloat162_rn(const float a)
Converts input to nv_bfloat16 precision in round-to-nearest-even mode and populates both
halves of nv_bfloat162 with converted value.
__host__ __device__ __nv_bfloat16 __float2bfloat16_rd(const float a)
Converts float number to nv_bfloat16 precision in round-down mode and returns nv_bfloat16
with converted value.
5.5. Bfloat16 Precision Conversion and Data Movement
159
__host__ __device__ __nv_bfloat16 __float2bfloat16_rn(const float a)
Converts float number to nv_bfloat16 precision in round-to-nearest-even mode and returns
nv_bfloat16 with converted value.
__host__ __device__ __nv_bfloat16 __float2bfloat16_ru(const float a)
Converts float number to nv_bfloat16 precision in round-up mode and returns nv_bfloat16
with converted value.
__host__ __device__ __nv_bfloat16 __float2bfloat16_rz(const float a)
Converts float number to nv_bfloat16 precision in round-towards-zero mode and returns
nv_bfloat16 with converted value.
__host__ __device__ __nv_bfloat162 __floats2bfloat162_rn(const float a, const float b)
Converts both input floats to nv_bfloat16 precision in round-to-nearest-even mode and returns
nv_bfloat162 with converted values.
__host__ __device__ __nv_bfloat162 __halves2bfloat162(const __nv_bfloat16 a, const __nv_bfloat16
b)
Combines two nv_bfloat16 numbers into one nv_bfloat162 number.
__host__ __device__ __nv_bfloat16 __high2bfloat16(const __nv_bfloat162 a)
Returns high 16 bits of nv_bfloat162 input.
__host__ __device__ __nv_bfloat162 __high2bfloat162(const __nv_bfloat162 a)
Extracts high 16 bits from nv_bfloat162 input.
__host__ __device__ float __high2float(const __nv_bfloat162 a)
Converts high 16 bits of nv_bfloat162 to float and returns the result.
__host__
__device__
__nv_bfloat162
__highs2bfloat162(const
__nv_bfloat162 a,
const
__nv_bfloat162 b)
Extracts high 16 bits from each of the two nv_bfloat162 inputs and combines into one
nv_bfloat162 number.
__device__ __nv_bfloat16 __int2bfloat16_rd(const int i)
Convert a signed integer to a nv_bfloat16 in round-down mode.
__host__ __device__ __nv_bfloat16 __int2bfloat16_rn(const int i)
Convert a signed integer to a nv_bfloat16 in round-to-nearest-even mode.
__device__ __nv_bfloat16 __int2bfloat16_ru(const int i)
Convert a signed integer to a nv_bfloat16 in round-up mode.
__device__ __nv_bfloat16 __int2bfloat16_rz(const int i)
Convert a signed integer to a nv_bfloat16 in round-towards-zero mode.
__device__ __nv_bfloat162 __ldca(const __nv_bfloat162 *const ptr)
Generates a ld.global.ca load instruction.
__device__ __nv_bfloat16 __ldca(const __nv_bfloat16 *const ptr)
Generates a ld.global.ca load instruction.
__device__ __nv_bfloat16 __ldcg(const __nv_bfloat16 *const ptr)
Generates a ld.global.cg load instruction.
__device__ __nv_bfloat162 __ldcg(const __nv_bfloat162 *const ptr)
Generates a ld.global.cg load instruction.
__device__ __nv_bfloat162 __ldcs(const __nv_bfloat162 *const ptr)
Generates a ld.global.cs load instruction.
__device__ __nv_bfloat16 __ldcs(const __nv_bfloat16 *const ptr)
Generates a ld.global.cs load instruction.
160
Chapter 5. Bfloat16 Precision Intrinsics
__device__ __nv_bfloat16 __ldcv(const __nv_bfloat16 *const ptr)
Generates a ld.global.cv load instruction.
__device__ __nv_bfloat162 __ldcv(const __nv_bfloat162 *const ptr)
Generates a ld.global.cv load instruction.
__device__ __nv_bfloat162 __ldg(const __nv_bfloat162 *const ptr)
Generates a ld.global.nc load instruction.
__device__ __nv_bfloat16 __ldg(const __nv_bfloat16 *const ptr)
Generates a ld.global.nc load instruction.
__device__ __nv_bfloat162 __ldlu(const __nv_bfloat162 *const ptr)
Generates a ld.global.lu load instruction.
__device__ __nv_bfloat16 __ldlu(const __nv_bfloat16 *const ptr)
Generates a ld.global.lu load instruction.
__device__ __nv_bfloat16 __ll2bfloat16_rd(const long long int i)
Convert a signed 64-bit integer to a nv_bfloat16 in round-down mode.
__host__ __device__ __nv_bfloat16 __ll2bfloat16_rn(const long long int i)
Convert a signed 64-bit integer to a nv_bfloat16 in round-to-nearest-even mode.
__device__ __nv_bfloat16 __ll2bfloat16_ru(const long long int i)
Convert a signed 64-bit integer to a nv_bfloat16 in round-up mode.
__device__ __nv_bfloat16 __ll2bfloat16_rz(const long long int i)
Convert a signed 64-bit integer to a nv_bfloat16 in round-towards-zero mode.
__host__ __device__ __nv_bfloat16 __low2bfloat16(const __nv_bfloat162 a)
Returns low 16 bits of nv_bfloat162 input.
__host__ __device__ __nv_bfloat162 __low2bfloat162(const __nv_bfloat162 a)
Extracts low 16 bits from nv_bfloat162 input.
__host__ __device__ float __low2float(const __nv_bfloat162 a)
Converts low 16 bits of nv_bfloat162 to float and returns the result.
__host__ __device__ __nv_bfloat162 __lowhigh2highlow(const __nv_bfloat162 a)
Swaps both halves of the nv_bfloat162 input.
__host__
__device__
__nv_bfloat162
__lows2bfloat162(const
__nv_bfloat162
a,
const
__nv_bfloat162 b)
Extracts low 16 bits from each of the two nv_bfloat162 inputs and combines into one
nv_bfloat162 number.
__nv_bfloat162::__nv_bfloat162()=default
Constructor by default.
__host__ __device__ constexpr __nv_bfloat162::__nv_bfloat162(const __nv_bfloat16 &a, const
__nv_bfloat16 &b)
Constructor from two __nv_bfloat16 variables.
__host__ __device__ __nv_bfloat162::__nv_bfloat162(__nv_bfloat162 &&src)
Move constructor, available for C++11 and later dialects.
__host__ __device__ __nv_bfloat162::__nv_bfloat162(const __nv_bfloat162_raw &h2r)
Constructor from __nv_bfloat162_raw .
__host__ __device__ __nv_bfloat162::__nv_bfloat162(const __nv_bfloat162 &src)
Copy constructor.
5.5. Bfloat16 Precision Conversion and Data Movement
161
__host__ __device__ __nv_bfloat162::operator __nv_bfloat162_raw() const
Conversion operator to __nv_bfloat162_raw .
__host__ __device__ __nv_bfloat162 & __nv_bfloat162::operator=(const __nv_bfloat162 &src)
Copy assignment operator.
__host__ __device__ __nv_bfloat162 & __nv_bfloat162::operator=(const __nv_bfloat162_raw &h2r)
Assignment operator from __nv_bfloat162_raw .
__host__ __device__ __nv_bfloat162 & __nv_bfloat162::operator=(__nv_bfloat162 &&src)
Move assignment operator, available for C++11 and later dialects.
__host__ __device__ __nv_bfloat16::__nv_bfloat16(const double f)
Construct __nv_bfloat16 from double input using default round-to-nearest-even rounding
mode.
__host__ __device__ __nv_bfloat16::__nv_bfloat16(const float f)
Construct __nv_bfloat16 from float input using default round-to-nearest-even rounding
mode.
__host__ __device__ __nv_bfloat16::__nv_bfloat16(long long val)
Construct __nv_bfloat16 from long long input using default round-to-nearest-even rounding
mode.
__host__ __device__ __nv_bfloat16::__nv_bfloat16(unsigned short val)
Construct __nv_bfloat16 from unsigned short integer input using default round-to-nearest-
even rounding mode.
__host__ __device__ __nv_bfloat16::__nv_bfloat16(unsigned int val)
Construct __nv_bfloat16 from unsigned int input using default round-to-nearest-even
rounding mode.
__nv_bfloat16::__nv_bfloat16()=default
Constructor by default.
__host__ __device__ __nv_bfloat16::__nv_bfloat16(const __half f)
Construct __nv_bfloat16 from __half input using default round-to-nearest-even rounding
mode.
__host__ __device__ __nv_bfloat16::__nv_bfloat16(short val)
Construct __nv_bfloat16 from short integer input using default round-to-nearest-even
rounding mode.
__host__ __device__ constexpr __nv_bfloat16::__nv_bfloat16(const __nv_bfloat16_raw &hr)
Constructor from __nv_bfloat16_raw .
__host__ __device__ __nv_bfloat16::__nv_bfloat16(unsigned long long val)
Construct __nv_bfloat16 from unsigned long long input using default round-to-nearest-
even rounding mode.
__host__ __device__ __nv_bfloat16::__nv_bfloat16(int val)
Construct __nv_bfloat16 from int input using default round-to-nearest-even rounding mode.
__host__ __device__ __nv_bfloat16::__nv_bfloat16(const unsigned long val)
Construct __nv_bfloat16 from unsigned long input using default round-to-nearest-even
rounding mode.
__host__ __device__ __nv_bfloat16::__nv_bfloat16(const long val)
Construct __nv_bfloat16 from long input using default round-to-nearest-even rounding
mode.
162
Chapter 5. Bfloat16 Precision Intrinsics
__host__ __device__ __nv_bfloat16::operator __nv_bfloat16_raw() const
Type cast to __nv_bfloat16_raw operator.
__host__ __device__ __nv_bfloat16::operator __nv_bfloat16_raw() const volatile
Type cast to __nv_bfloat16_raw operator with volatile input.
__host__ __device__ constexpr __nv_bfloat16::operator bool() const
Conversion operator to bool data type.
__host__ __device__ __nv_bfloat16::operator char() const
Conversion operator to an implementation defined char data type.
__host__ __device__ __nv_bfloat16::operator float() const
Type cast to float operator.
__host__ __device__ __nv_bfloat16::operator int() const
Conversion operator to int data type.
__host__ __device__ __nv_bfloat16::operator long() const
Conversion operator to long data type.
__host__ __device__ __nv_bfloat16::operator long long() const
Conversion operator to long long data type.
__host__ __device__ __nv_bfloat16::operator short() const
Conversion operator to short data type.
__host__ __device__ __nv_bfloat16::operator signed char() const
Conversion operator to signed char data type.
__host__ __device__ __nv_bfloat16::operator unsigned char() const
Conversion operator to unsigned char data type.
__host__ __device__ __nv_bfloat16::operator unsigned int() const
Conversion operator to unsigned int data type.
__host__ __device__ __nv_bfloat16::operator unsigned long() const
Conversion operator to unsigned long data type.
__host__ __device__ __nv_bfloat16::operator unsigned long long() const
Conversion operator to unsigned long long data type.
__host__ __device__ __nv_bfloat16::operator unsigned short() const
Conversion operator to unsigned short data type.
__host__ __device__ __nv_bfloat16 & __nv_bfloat16::operator=(const __nv_bfloat16_raw &hr)
Assignment operator from __nv_bfloat16_raw .
__host__ __device__ __nv_bfloat16 & __nv_bfloat16::operator=(unsigned int val)
Type cast from unsigned int assignment operator, using default round-to-nearest-even round-
ing mode.
__host__ __device__ __nv_bfloat16 & __nv_bfloat16::operator=(int val)
Type cast from int assignment operator, using default round-to-nearest-even rounding mode.
__host__ __device__ __nv_bfloat16 & __nv_bfloat16::operator=(long long val)
Type cast from long long assignment operator, using default round-to-nearest-even rounding
mode.
__host__ __device__ __nv_bfloat16 & __nv_bfloat16::operator=(unsigned long long val)
Type cast from unsigned long long assignment operator, using default round-to-nearest-even
rounding mode.
5.5. Bfloat16 Precision Conversion and Data Movement
163
__host__ __device__ __nv_bfloat16 & __nv_bfloat16::operator=(unsigned short val)
Type cast from unsigned short assignment operator, using default round-to-nearest-even
rounding mode.
__host__ __device__ volatile __nv_bfloat16 & __nv_bfloat16::operator=(const __nv_bfloat16_raw
&hr) volatile
Assignment operator from __nv_bfloat16_raw to volatile __nv_bfloat16 .
__host__ __device__ __nv_bfloat16 & __nv_bfloat16::operator=(const double f)
Type cast to __nv_bfloat16 assignment operator from double input using default round-to-
nearest-even rounding mode.
__host__ __device__ __nv_bfloat16 & __nv_bfloat16::operator=(const float f)
Type cast to __nv_bfloat16 assignment operator from float input using default round-to-
nearest-even rounding mode.
__host__
__device__
volatile
__nv_bfloat16
&
__nv_bfloat16::operator=(const
volatile
__nv_bfloat16_raw &hr) volatile
Assignment operator from volatile __nv_bfloat16_raw to volatile __nv_bfloat16 .
__host__ __device__ __nv_bfloat16 & __nv_bfloat16::operator=(short val)
Type cast from short assignment operator, using default round-to-nearest-even rounding mode.
__device__ __nv_bfloat162 __shfl_down_sync(const unsigned int mask, const __nv_bfloat162 var,
const unsigned int delta, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __nv_bfloat16 __shfl_down_sync(const unsigned int mask, const __nv_bfloat16 var,
const unsigned int delta, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __nv_bfloat162 __shfl_sync(const unsigned int mask, const __nv_bfloat162 var, const int
srcLane, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __nv_bfloat16 __shfl_sync(const unsigned int mask, const __nv_bfloat16 var, const int
srcLane, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __nv_bfloat16 __shfl_up_sync(const unsigned int mask, const __nv_bfloat16 var, const
unsigned int delta, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __nv_bfloat162 __shfl_up_sync(const unsigned int mask, const __nv_bfloat162 var, const
unsigned int delta, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __nv_bfloat16 __shfl_xor_sync(const unsigned int mask, const __nv_bfloat16 var, const
int laneMask, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __nv_bfloat162 __shfl_xor_sync(const unsigned int mask, const __nv_bfloat162 var,
const int laneMask, const int width=warpSize)
Exchange a variable between threads within a warp.
__device__ __nv_bfloat16 __short2bfloat16_rd(const short int i)
Convert a signed short integer to a nv_bfloat16 in round-down mode.
__host__ __device__ __nv_bfloat16 __short2bfloat16_rn(const short int i)
Convert a signed short integer to a nv_bfloat16 in round-to-nearest-even mode.
164
Chapter 5. Bfloat16 Precision Intrinsics
__device__ __nv_bfloat16 __short2bfloat16_ru(const short int i)
Convert a signed short integer to a nv_bfloat16 in round-up mode.
__device__ __nv_bfloat16 __short2bfloat16_rz(const short int i)
Convert a signed short integer to a nv_bfloat16 in round-towards-zero mode.
__host__ __device__ __nv_bfloat16 __short_as_bfloat16(const short int i)
Reinterprets bits in a signed short integer as a nv_bfloat16 .
__device__ void __stcg(__nv_bfloat16 *const ptr, const __nv_bfloat16 value)
Generates a st.global.cg store instruction.
__device__ void __stcg(__nv_bfloat162 *const ptr, const __nv_bfloat162 value)
Generates a st.global.cg store instruction.
__device__ void __stcs(__nv_bfloat16 *const ptr, const __nv_bfloat16 value)
Generates a st.global.cs store instruction.
__device__ void __stcs(__nv_bfloat162 *const ptr, const __nv_bfloat162 value)
Generates a st.global.cs store instruction.
__device__ void __stwb(__nv_bfloat16 *const ptr, const __nv_bfloat16 value)
Generates a st.global.wb store instruction.
__device__ void __stwb(__nv_bfloat162 *const ptr, const __nv_bfloat162 value)
Generates a st.global.wb store instruction.
__device__ void __stwt(__nv_bfloat162 *const ptr, const __nv_bfloat162 value)
Generates a st.global.wt store instruction.
__device__ void __stwt(__nv_bfloat16 *const ptr, const __nv_bfloat16 value)
Generates a st.global.wt store instruction.
__device__ __nv_bfloat16 __uint2bfloat16_rd(const unsigned int i)
Convert an unsigned integer to a nv_bfloat16 in round-down mode.
__host__ __device__ __nv_bfloat16 __uint2bfloat16_rn(const unsigned int i)
Convert an unsigned integer to a nv_bfloat16 in round-to-nearest-even mode.
__device__ __nv_bfloat16 __uint2bfloat16_ru(const unsigned int i)
Convert an unsigned integer to a nv_bfloat16 in round-up mode.
__device__ __nv_bfloat16 __uint2bfloat16_rz(const unsigned int i)
Convert an unsigned integer to a nv_bfloat16 in round-towards-zero mode.
__device__ __nv_bfloat16 __ull2bfloat16_rd(const unsigned long long int i)
Convert an unsigned 64-bit integer to a nv_bfloat16 in round-down mode.
__host__ __device__ __nv_bfloat16 __ull2bfloat16_rn(const unsigned long long int i)
Convert an unsigned 64-bit integer to a nv_bfloat16 in round-to-nearest-even mode.
__device__ __nv_bfloat16 __ull2bfloat16_ru(const unsigned long long int i)
Convert an unsigned 64-bit integer to a nv_bfloat16 in round-up mode.
__device__ __nv_bfloat16 __ull2bfloat16_rz(const unsigned long long int i)
Convert an unsigned 64-bit integer to a nv_bfloat16 in round-towards-zero mode.
__device__ __nv_bfloat16 __ushort2bfloat16_rd(const unsigned short int i)
Convert an unsigned short integer to a nv_bfloat16 in round-down mode.
__host__ __device__ __nv_bfloat16 __ushort2bfloat16_rn(const unsigned short int i)
Convert an unsigned short integer to a nv_bfloat16 in round-to-nearest-even mode.
5.5. Bfloat16 Precision Conversion and Data Movement
165
__device__ __nv_bfloat16 __ushort2bfloat16_ru(const unsigned short int i)
Convert an unsigned short integer to a nv_bfloat16 in round-up mode.
__device__ __nv_bfloat16 __ushort2bfloat16_rz(const unsigned short int i)
Convert an unsigned short integer to a nv_bfloat16 in round-towards-zero mode.
__host__ __device__ __nv_bfloat16 __ushort_as_bfloat16(const unsigned short int i)
Reinterprets bits in an unsigned short integer as a nv_bfloat16 .
__host__ __device__ __nv_bfloat162 make_bfloat162(const __nv_bfloat16 x, const __nv_bfloat16 y)
Vector function, combines two nv_bfloat16 numbers into one nv_bfloat162 number.
5.5.1. Functions
__host__ __device__ float2 __bfloat1622float2(const __nv_bfloat162 a)
Converts both halves of nv_bfloat162 to float2 and returns the result.
Converts both halves of nv_bfloat162 input a to float and returns the result as a float2 packed
value.
See also:
__bfloat162float(__nv_bfloat16) for further details.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
float2
▶ a converted to float2.
__host__ __device__ __nv_bfloat162 __bfloat162bfloat162(const __nv_bfloat16 a)
Returns nv_bfloat162 with both halves equal to the input value.
Returns nv_bfloat162 number with both halves equal to the input a nv_bfloat16 number.
Parameters
a - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat162
▶ The vector which has both its halves equal to the input a.
__host__ __device__ signed char __bfloat162char_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed char in round-towards-zero mode.
Convert the nv_bfloat16 floating-point value h to a signed char in round-towards-zero mode. NaN
inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
signed char
▶ h converted to a signed char using round-towards-zero mode.
▶ __bfloat162char_rz (±0) returns 0.
166
Chapter 5. Bfloat16 Precision Intrinsics
▶ __bfloat162char_rz (x), x > 127 returns SCHAR_MAX = 0x7F.
▶ __bfloat162char_rz (x), x < -128 returns SCHAR_MIN = 0x80.
▶ __bfloat162char_rz(NaN) returns 0.
__host__ __device__ float __bfloat162float(const __nv_bfloat16 a)
Converts nv_bfloat16 number to float.
Converts nv_bfloat16 number a to float.
Parameters
a - [in] - float. Is only being read.
Returns
float
▶ a converted to float.
▶ __bfloat162float (±0) returns ±0.
▶ __bfloat162float (±∞) returns ±∞.
▶ __bfloat162float(NaN) returns NaN.
__device__ int __bfloat162int_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed integer in round-down mode.
Convert the nv_bfloat16 floating-point value h to a signed integer in round-down mode. NaN
inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
int
▶ h converted to a signed integer using round-down mode.
▶ __bfloat162int_rd (±0) returns 0.
▶ __bfloat162int_rd (x), x > INTM AX returns INT_MAX = 0x7FFFFFFF.
▶ __bfloat162int_rd (x), x < INTM IN returns INT_MIN = 0x80000000.
▶ __bfloat162int_rd(NaN) returns 0.*
__device__ int __bfloat162int_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed integer in round-to-nearest-even mode.
Convert the nv_bfloat16 floating-point value h to a signed integer in round-to-nearest-even
mode. NaN inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
int
▶ h converted to a signed integer using round-to-nearest-even mode.
▶ __bfloat162int_rn (±0) returns 0.
▶ __bfloat162int_rn (x), x > INTM AX returns INT_MAX = 0x7FFFFFFF.
▶ __bfloat162int_rn (x), x < INTM IN returns INT_MIN = 0x80000000.
5.5. Bfloat16 Precision Conversion and Data Movement
167
▶ __bfloat162int_rn(NaN) returns 0.
__device__ int __bfloat162int_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed integer in round-up mode.
Convert the nv_bfloat16 floating-point value h to a signed integer in round-up mode. NaN inputs
are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
int
▶ h converted to a signed integer using round-up mode.
▶ __bfloat162int_ru (±0) returns 0.
▶ __bfloat162int_ru (x), x > INTM AX returns INT_MAX = 0x7FFFFFFF.
▶ __bfloat162int_ru (x), x < INTM IN returns INT_MIN = 0x80000000.
▶ __bfloat162int_ru(NaN) returns 0.
__host__ __device__ int __bfloat162int_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed integer in round-towards-zero mode.
Convert the nv_bfloat16 floating-point value h to a signed integer in round-towards-zero mode.
NaN inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
int
▶ h converted to a signed integer using round-towards-zero mode.
▶ __bfloat162int_rz (±0) returns 0.
▶ __bfloat162int_rz (x), x > INTM AX returns INT_MAX = 0x7FFFFFFF.
▶ __bfloat162int_rz (x), x < INTM IN returns INT_MIN = 0x80000000.
▶ __bfloat162int_rz(NaN) returns 0.
__device__ long long int __bfloat162ll_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed 64-bit integer in round-down mode.
Convert the nv_bfloat16 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] - nv_bfloat16. Is only being read.
Returns
long long int
▶ h converted to a signed 64-bit integer.
__device__ long long int __bfloat162ll_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed 64-bit integer in round-to-nearest-even mode.
Convert the nv_bfloat16 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.
168
Chapter 5. Bfloat16 Precision Intrinsics
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
long long int
▶ h converted to a signed 64-bit integer.
__device__ long long int __bfloat162ll_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed 64-bit integer in round-up mode.
Convert the nv_bfloat16 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] - nv_bfloat16. Is only being read.
Returns
long long int
▶ h converted to a signed 64-bit integer.
__host__ __device__ long long int __bfloat162ll_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed 64-bit integer in round-towards-zero mode.
Convert the nv_bfloat16 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] - nv_bfloat16. Is only being read.
Returns
long long int
▶ h converted to a signed 64-bit integer.
__device__ short int __bfloat162short_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed short integer in round-down mode.
Convert the nv_bfloat16 floating-point value h to a signed short integer in round-down mode.
NaN inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
short int
▶ h converted to a signed short integer using round-down mode.
▶ __bfloat162short_rd (±0) returns 0.
▶ __bfloat162short_rd (x), x > 32767 returns SHRT_MAX = 0x7FFF.
▶ __bfloat162short_rd (x), x < -32768 returns SHRT_MIN = 0x8000.
▶ __bfloat162short_rd(NaN) returns 0.
__device__ short int __bfloat162short_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed short integer in round-to-nearest-even mode.
Convert the nv_bfloat16 floating-point value h to a signed short integer in round-to-nearest-even
mode. NaN inputs are converted to 0.
5.5. Bfloat16 Precision Conversion and Data Movement
169
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
short int
▶ h converted to a signed short integer using round-to-nearest-even mode.
▶ __bfloat162short_rn (±0) returns 0.
▶ __bfloat162short_rn (x), x > 32767 returns SHRT_MAX = 0x7FFF.
▶ __bfloat162short_rn (x), x < -32768 returns SHRT_MIN = 0x8000.
▶ __bfloat162short_rn(NaN) returns 0.
__device__ short int __bfloat162short_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed short integer in round-up mode.
Convert the nv_bfloat16 floating-point value h to a signed short integer in round-up mode. NaN
inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
short int
▶ h converted to a signed short integer using round-up mode.
▶ __bfloat162short_ru (±0) returns 0.
▶ __bfloat162short_ru (x), x > 32767 returns SHRT_MAX = 0x7FFF.
▶ __bfloat162short_ru (x), x < -32768 returns SHRT_MIN = 0x8000.
▶ __bfloat162short_ru(NaN) returns 0.
__host__ __device__ short int __bfloat162short_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to a signed short integer in round-towards-zero mode.
Convert the nv_bfloat16 floating-point value h to a signed short integer in round-towards-zero
mode. NaN inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
short int
▶ h converted to a signed short integer using round-towards-zero mode.
▶ __bfloat162short_rz (±0) returns 0.
▶ __bfloat162short_rz (x), x > 32767 returns SHRT_MAX = 0x7FFF.
▶ __bfloat162short_rz (x), x < -32768 returns SHRT_MIN = 0x8000.
▶ __bfloat162short_rz(NaN) returns 0.
__host__ __device__ unsigned char __bfloat162uchar_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned char in round-towards-zero mode.
Convert the nv_bfloat16 floating-point value h to an unsigned char in round-towards-zero mode.
NaN inputs are converted to 0.
170
Chapter 5. Bfloat16 Precision Intrinsics
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned char
▶ h converted to an unsigned char using round-towards-zero mode.
▶ __bfloat162uchar_rz (±0) returns 0.
▶ __bfloat162uchar_rz (x), x > 255 returns UCHAR_MAX = 0xFF.
▶ __bfloat162uchar_rz (x), x < 0.0 returns 0.
▶ __bfloat162uchar_rz(NaN) returns 0.
__device__ unsigned int __bfloat162uint_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned integer in round-down mode.
Convert the nv_bfloat16 floating-point value h to an unsigned integer in round-down mode. NaN
inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned int
▶ h converted to an unsigned integer.
__device__ unsigned int __bfloat162uint_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned integer in round-to-nearest-even mode.
Convert the nv_bfloat16 floating-point value h to an unsigned integer in round-to-nearest-even
mode. NaN inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned int
▶ h converted to an unsigned integer.
__device__ unsigned int __bfloat162uint_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned integer in round-up mode.
Convert the nv_bfloat16 floating-point value h to an unsigned integer in round-up mode. NaN
inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned int
▶ h converted to an unsigned integer.
__host__ __device__ unsigned int __bfloat162uint_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned integer in round-towards-zero mode.
Convert the nv_bfloat16 floating-point value h to an unsigned integer in round-towards-zero
mode. NaN inputs are converted to 0.
5.5. Bfloat16 Precision Conversion and Data Movement
171
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned int
▶ h converted to an unsigned integer.
__device__ unsigned long long int __bfloat162ull_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned 64-bit integer in round-down mode.
Convert the nv_bfloat16 floating-point value h to an unsigned 64-bit integer in round-down mode.
NaN inputs return 0x8000000000000000.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned long long int
▶ h converted to an unsigned 64-bit integer.
__device__ unsigned long long int __bfloat162ull_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned 64-bit integer in round-to-nearest-even mode.
Convert the nv_bfloat16 floating-point value h to an unsigned 64-bit integer in round-to-nearest-
even mode. NaN inputs return 0x8000000000000000.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned long long int
▶ h converted to an unsigned 64-bit integer.
__device__ unsigned long long int __bfloat162ull_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned 64-bit integer in round-up mode.
Convert the nv_bfloat16 floating-point value h to an unsigned 64-bit integer in round-up mode.
NaN inputs return 0x8000000000000000.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned long long int
▶ h converted to an unsigned 64-bit integer.
__host__ __device__ unsigned long long int __bfloat162ull_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned 64-bit integer in round-towards-zero mode.
Convert the nv_bfloat16 floating-point value h to an unsigned 64-bit integer in round-towards-
zero mode. NaN inputs return 0x8000000000000000.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned long long int
▶ h converted to an unsigned 64-bit integer.
172
Chapter 5. Bfloat16 Precision Intrinsics
__device__ unsigned short int __bfloat162ushort_rd(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned short integer in round-down mode.
Convert the nv_bfloat16 floating-point value h to an unsigned short integer in round-down mode.
NaN inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned short int
▶ h converted to an unsigned short integer.
__device__ unsigned short int __bfloat162ushort_rn(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned short integer in round-to-nearest-even mode.
Convert the nv_bfloat16 floating-point value h to an unsigned short integer in round-to-nearest-
even mode. NaN inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned short int
▶ h converted to an unsigned short integer.
__device__ unsigned short int __bfloat162ushort_ru(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned short integer in round-up mode.
Convert the nv_bfloat16 floating-point value h to an unsigned short integer in round-up mode.
NaN inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned short int
▶ h converted to an unsigned short integer.
__host__ __device__ unsigned short int __bfloat162ushort_rz(const __nv_bfloat16 h)
Convert a nv_bfloat16 to an unsigned short integer in round-towards-zero mode.
Convert the nv_bfloat16 floating-point value h to an unsigned short integer in round-towards-
zero mode. NaN inputs are converted to 0.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned short int
▶ h converted to an unsigned short integer.
__host__ __device__ short int __bfloat16_as_short(const __nv_bfloat16 h)
Reinterprets bits in a nv_bfloat16 as a signed short integer.
Reinterprets the bits in the nv_bfloat16 floating-point number h as a signed short integer.
Parameters
h - [in] - nv_bfloat16. Is only being read.
5.5. Bfloat16 Precision Conversion and Data Movement
173
Returns
short int
▶ The reinterpreted value.
__host__ __device__ unsigned short int __bfloat16_as_ushort(const __nv_bfloat16 h)
Reinterprets bits in a nv_bfloat16 as an unsigned short integer.
Reinterprets the bits in the nv_bfloat16 floating-point h as an unsigned short number.
Parameters
h - [in] - nv_bfloat16. Is only being read.
Returns
unsigned short int
▶ The reinterpreted value.
__host__ __device__ __nv_bfloat16 __double2bfloat16(const double a)
Converts double number to nv_bfloat16 precision in round-to-nearest-even mode and returns
nv_bfloat16 with converted value.
Converts double number a to nv_bfloat16 precision in round-to-nearest-even mode.
Parameters
a - [in] - double. Is only being read.
Returns
nv_bfloat16
▶ a converted to nv_bfloat16 using round-to-nearest-even mode.
▶ __double2bfloat16 (±0) returns ±0.
▶ __double2bfloat16 (±∞) returns ±∞.
▶ __double2bfloat16(NaN) returns NaN.
__host__ __device__ __nv_bfloat162 __float22bfloat162_rn(const float2 a)
Converts both components of float2 number to nv_bfloat16 precision in round-to-nearest-even
mode and returns nv_bfloat162 with converted values.
Converts both components of float2 to nv_bfloat16 precision in round-to-nearest-even mode
and combines the results into one nv_bfloat162 number. Low 16 bits of the return value cor-
respond to a.x and high 16 bits of the return value correspond to a.y.
See also:
__float2bfloat16_rn(float) for further details.
Parameters
a - [in] - float2. Is only being read.
Returns
nv_bfloat162
▶ The nv_bfloat162 which has corresponding halves equal to the converted
float2 components.
__host__ __device__ __nv_bfloat16 __float2bfloat16(const float a)
Converts float number to nv_bfloat16 precision in round-to-nearest-even mode and returns
nv_bfloat16 with converted value.
174
Chapter 5. Bfloat16 Precision Intrinsics
Converts float number a to nv_bfloat16 precision in round-to-nearest-even mode.
See also:
__float2bfloat16_rn(float) for further details.
Parameters
a - [in] - float. Is only being read.
Returns
nv_bfloat16
▶ a converted to nv_bfloat16 using round-to-nearest-even mode.
__host__ __device__ __nv_bfloat162 __float2bfloat162_rn(const float a)
Converts input to nv_bfloat16 precision in round-to-nearest-even mode and populates both
halves of nv_bfloat162 with converted value.
Converts input a to nv_bfloat16 precision in round-to-nearest-even mode and populates both
halves of nv_bfloat162 with converted value.
See also:
__float2bfloat16_rn(float) for further details.
Parameters
a - [in] - float. Is only being read.
Returns
nv_bfloat162
▶ The nv_bfloat162 value with both halves equal to the converted nv_bfloat16
precision number.
__host__ __device__ __nv_bfloat16 __float2bfloat16_rd(const float a)
Converts float number to nv_bfloat16 precision in round-down mode and returns nv_bfloat16
with converted value.
Converts float number a to nv_bfloat16 precision in round-down mode.
Parameters
a - [in] - float. Is only being read.
Returns
nv_bfloat16
▶ a converted to nv_bfloat16 using round-down mode.
▶ __float2bfloat16_rd (±0) returns ±0.
▶ __float2bfloat16_rd (±∞) returns ±∞.
▶ __float2bfloat16_rd(NaN) returns NaN.
__host__ __device__ __nv_bfloat16 __float2bfloat16_rn(const float a)
Converts float number to nv_bfloat16 precision in round-to-nearest-even mode and returns
nv_bfloat16 with converted value.
Converts float number a to nv_bfloat16 precision in round-to-nearest-even mode.
Parameters
a - [in] - float. Is only being read.
5.5. Bfloat16 Precision Conversion and Data Movement
175
Returns
nv_bfloat16
▶ a converted to nv_bfloat16 using round-to-nearest-even mode.
▶ __float2bfloat16_rn (±0) returns ±0.
▶ __float2bfloat16_rn (±∞) returns ±∞.
▶ __float2bfloat16_rn(NaN) returns NaN.
__host__ __device__ __nv_bfloat16 __float2bfloat16_ru(const float a)
Converts float number to nv_bfloat16 precision in round-up mode and returns nv_bfloat16
with converted value.
Converts float number a to nv_bfloat16 precision in round-up mode.
Parameters
a - [in] - float. Is only being read.
Returns
nv_bfloat16
▶ a converted to nv_bfloat16 using round-up mode.
▶ __float2bfloat16_ru (±0) returns ±0.
▶ __float2bfloat16_ru (±∞) returns ±∞.
▶ __float2bfloat16_ru(NaN) returns NaN.
__host__ __device__ __nv_bfloat16 __float2bfloat16_rz(const float a)
Converts float number to nv_bfloat16 precision in round-towards-zero mode and returns
nv_bfloat16 with converted value.
Converts float number a to nv_bfloat16 precision in round-towards-zero mode.
Parameters
a - [in] - float. Is only being read.
Returns
nv_bfloat16
▶ a converted to nv_bfloat16 using round-towards-zero mode.
▶ __float2bfloat16_rz (±0) returns ±0.
▶ __float2bfloat16_rz (±∞) returns ±∞.
▶ __float2bfloat16_rz(NaN) returns NaN.
__host__ __device__ __nv_bfloat162 __floats2bfloat162_rn(const float a, const float b)
Converts both input floats to nv_bfloat16 precision in round-to-nearest-even mode and returns
nv_bfloat162 with converted values.
Converts both input floats to nv_bfloat16 precision in round-to-nearest-even mode and com-
bines the results into one nv_bfloat162 number. Low 16 bits of the return value correspond to
the input a, high 16 bits correspond to the input b.
See also:
__float2bfloat16_rn(float) for further details.
Parameters
176
Chapter 5. Bfloat16 Precision Intrinsics
▶ a - [in] - float. Is only being read.
▶ b - [in] - float. Is only being read.
Returns
nv_bfloat162
▶ The nv_bfloat162 value with corresponding halves equal to the converted in-
put floats.
__host__ __device__ __nv_bfloat162 __halves2bfloat162(const __nv_bfloat16 a, const
__nv_bfloat16 b)
Combines two nv_bfloat16 numbers into one nv_bfloat162 number.
Combines two input nv_bfloat16 number a and b into one nv_bfloat162 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.
Parameters
▶ a - [in] - nv_bfloat16. Is only being read.
▶ b - [in] - nv_bfloat16. Is only being read.
Returns
nv_bfloat162
▶ The nv_bfloat162 with one nv_bfloat16 equal to a and the other to b.
__host__ __device__ __nv_bfloat16 __high2bfloat16(const __nv_bfloat162 a)
Returns high 16 bits of nv_bfloat162 input.
Returns high 16 bits of nv_bfloat162 input a.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat16
▶ The high 16 bits of the input.
__host__ __device__ __nv_bfloat162 __high2bfloat162(const __nv_bfloat162 a)
Extracts high 16 bits from nv_bfloat162 input.
Extracts high 16 bits from nv_bfloat162 input a and returns a new nv_bfloat162 number
which has both halves equal to the extracted bits.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The nv_bfloat162 with both halves equal to the high 16 bits of the input.
__host__ __device__ float __high2float(const __nv_bfloat162 a)
Converts high 16 bits of nv_bfloat162 to float and returns the result.
Converts high 16 bits of nv_bfloat162 input a to 32-bit floating-point number and returns the
result.
See also:
__bfloat162float(__nv_bfloat16) for further details.
5.5. Bfloat16 Precision Conversion and Data Movement
177
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
float
▶ The high 16 bits of a converted to float.
__host__ __device__ __nv_bfloat162 __highs2bfloat162(const __nv_bfloat162 a, const
__nv_bfloat162 b)
Extracts high 16 bits from each of the two nv_bfloat162 inputs and combines into one
nv_bfloat162 number.
Extracts high 16 bits from each of the two nv_bfloat162 inputs and combines into one
nv_bfloat162 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.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The high 16 bits of a and of b.
__device__ __nv_bfloat16 __int2bfloat16_rd(const int i)
Convert a signed integer to a nv_bfloat16 in round-down mode.
Convert the signed integer value i to a nv_bfloat16 floating-point value in round-down mode.
Parameters
i - [in] - int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__host__ __device__ __nv_bfloat16 __int2bfloat16_rn(const int i)
Convert a signed integer to a nv_bfloat16 in round-to-nearest-even mode.
Convert the signed integer value i to a nv_bfloat16 floating-point value in round-to-nearest-even
mode.
Parameters
i - [in] - int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __int2bfloat16_ru(const int i)
Convert a signed integer to a nv_bfloat16 in round-up mode.
Convert the signed integer value i to a nv_bfloat16 floating-point value in round-up mode.
Parameters
i - [in] - int. Is only being read.
178
Chapter 5. Bfloat16 Precision Intrinsics
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __int2bfloat16_rz(const int i)
Convert a signed integer to a nv_bfloat16 in round-towards-zero mode.
Convert the signed integer value i to a nv_bfloat16 floating-point value in round-towards-zero
mode.
Parameters
i - [in] - int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat162 __ldca(const __nv_bfloat162 *const ptr)
Generates a ld.global.ca load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __nv_bfloat16 __ldca(const __nv_bfloat16 *const ptr)
Generates a ld.global.ca load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __nv_bfloat16 __ldcg(const __nv_bfloat16 *const ptr)
Generates a ld.global.cg load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __nv_bfloat162 __ldcg(const __nv_bfloat162 *const ptr)
Generates a ld.global.cg load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __nv_bfloat162 __ldcs(const __nv_bfloat162 *const ptr)
Generates a ld.global.cs load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
5.5. Bfloat16 Precision Conversion and Data Movement
179
__device__ __nv_bfloat16 __ldcs(const __nv_bfloat16 *const ptr)
Generates a ld.global.cs load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __nv_bfloat16 __ldcv(const __nv_bfloat16 *const ptr)
Generates a ld.global.cv load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __nv_bfloat162 __ldcv(const __nv_bfloat162 *const ptr)
Generates a ld.global.cv load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __nv_bfloat162 __ldg(const __nv_bfloat162 *const ptr)
Generates a ld.global.nc load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __nv_bfloat16 __ldg(const __nv_bfloat16 *const ptr)
Generates a ld.global.nc load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __nv_bfloat162 __ldlu(const __nv_bfloat162 *const ptr)
Generates a ld.global.lu load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
__device__ __nv_bfloat16 __ldlu(const __nv_bfloat16 *const ptr)
Generates a ld.global.lu load instruction.
Parameters
ptr - [in] - memory location
Returns
The value pointed by ptr
180
Chapter 5. Bfloat16 Precision Intrinsics
__device__ __nv_bfloat16 __ll2bfloat16_rd(const long long int i)
Convert a signed 64-bit integer to a nv_bfloat16 in round-down mode.
Convert the signed 64-bit integer value i to a nv_bfloat16 floating-point value in round-down
mode.
Parameters
i - [in] - long long int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__host__ __device__ __nv_bfloat16 __ll2bfloat16_rn(const long long int i)
Convert a signed 64-bit integer to a nv_bfloat16 in round-to-nearest-even mode.
Convert the signed 64-bit integer value i to a nv_bfloat16 floating-point value in round-to-
nearest-even mode.
Parameters
i - [in] - long long int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __ll2bfloat16_ru(const long long int i)
Convert a signed 64-bit integer to a nv_bfloat16 in round-up mode.
Convert the signed 64-bit integer value i to a nv_bfloat16 floating-point value in round-up mode.
Parameters
i - [in] - long long int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __ll2bfloat16_rz(const long long int i)
Convert a signed 64-bit integer to a nv_bfloat16 in round-towards-zero mode.
Convert the signed 64-bit integer value i to a nv_bfloat16 floating-point value in round-towards-
zero mode.
Parameters
i - [in] - long long int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__host__ __device__ __nv_bfloat16 __low2bfloat16(const __nv_bfloat162 a)
Returns low 16 bits of nv_bfloat162 input.
Returns low 16 bits of nv_bfloat162 input a.
Parameters
a - [in] - nv_bfloat162. Is only being read.
5.5. Bfloat16 Precision Conversion and Data Movement
181
Returns
nv_bfloat16
▶ Returns nv_bfloat16 which contains low 16 bits of the input a.
__host__ __device__ __nv_bfloat162 __low2bfloat162(const __nv_bfloat162 a)
Extracts low 16 bits from nv_bfloat162 input.
Extracts low 16 bits from nv_bfloat162 input a and returns a new nv_bfloat162 number
which has both halves equal to the extracted bits.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The nv_bfloat162 with both halves equal to the low 16 bits of the input.
__host__ __device__ float __low2float(const __nv_bfloat162 a)
Converts low 16 bits of nv_bfloat162 to float and returns the result.
Converts low 16 bits of nv_bfloat162 input a to 32-bit floating-point number and returns the
result.
See also:
__bfloat162float(__nv_bfloat16) for further details.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
float
▶ The low 16 bits of a converted to float.
__host__ __device__ __nv_bfloat162 __lowhigh2highlow(const __nv_bfloat162 a)
Swaps both halves of the nv_bfloat162 input.
Swaps both halves of the nv_bfloat162 input and returns a new nv_bfloat162 number with
swapped halves.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ a with its halves being swapped.
__host__ __device__ __nv_bfloat162 __lows2bfloat162(const __nv_bfloat162 a, const
__nv_bfloat162 b)
Extracts low 16 bits from each of the two nv_bfloat162 inputs and combines into one
nv_bfloat162 number.
Extracts low 16 bits from each of the two nv_bfloat162 inputs and combines into one
nv_bfloat162 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
182
Chapter 5. Bfloat16 Precision Intrinsics
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The low 16 bits of a and of b.
__device__ __nv_bfloat162 __shfl_down_sync(const unsigned int mask, const __nv_bfloat162 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.
▶ 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] - nv_bfloat162. 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
nv_bfloat162.
__device__ __nv_bfloat16 __shfl_down_sync(const unsigned int mask, const __nv_bfloat16 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.
5.5. Bfloat16 Precision Conversion and Data Movement
183
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.
▶ 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] - nv_bfloat16. 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
nv_bfloat16.
__device__ __nv_bfloat162 __shfl_sync(const unsigned int mask, const __nv_bfloat162 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.
184
Chapter 5. Bfloat16 Precision Intrinsics
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] - nv_bfloat162. 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
nv_bfloat162.
__device__ __nv_bfloat16 __shfl_sync(const unsigned int mask, const __nv_bfloat16 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] - nv_bfloat16. Is only being read.
5.5. Bfloat16 Precision Conversion and Data Movement
185
▶ 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
nv_bfloat16.
__device__ __nv_bfloat16 __shfl_up_sync(const unsigned int mask, const __nv_bfloat16 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.
▶ 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] - nv_bfloat16. 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
nv_bfloat16.
__device__ __nv_bfloat162 __shfl_up_sync(const unsigned int mask, const __nv_bfloat162 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
186
Chapter 5. Bfloat16 Precision Intrinsics
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.
▶ 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] - nv_bfloat162. 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
nv_bfloat162.
__device__ __nv_bfloat16 __shfl_xor_sync(const unsigned int mask, const __nv_bfloat16 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
5.5. Bfloat16 Precision Conversion and Data Movement
187
▶ 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] - nv_bfloat16. 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
nv_bfloat16.
__device__ __nv_bfloat162 __shfl_xor_sync(const unsigned int mask, const __nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ laneMask - [in] - int. Is only being read.
▶ width - [in] - int. Is only being read.
188
Chapter 5. Bfloat16 Precision Intrinsics
Returns
Returns the 4-byte word referenced by var from the source thread ID as
nv_bfloat162.
__device__ __nv_bfloat16 __short2bfloat16_rd(const short int i)
Convert a signed short integer to a nv_bfloat16 in round-down mode.
Convert the signed short integer value i to a nv_bfloat16 floating-point value in round-down
mode.
Parameters
i - [in] - short int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__host__ __device__ __nv_bfloat16 __short2bfloat16_rn(const short int i)
Convert a signed short integer to a nv_bfloat16 in round-to-nearest-even mode.
Convert the signed short integer value i to a nv_bfloat16 floating-point value in round-to-
nearest-even mode.
Parameters
i - [in] - short int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __short2bfloat16_ru(const short int i)
Convert a signed short integer to a nv_bfloat16 in round-up mode.
Convert the signed short integer value i to a nv_bfloat16 floating-point value in round-up mode.
Parameters
i - [in] - short int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __short2bfloat16_rz(const short int i)
Convert a signed short integer to a nv_bfloat16 in round-towards-zero mode.
Convert the signed short integer value i to a nv_bfloat16 floating-point value in round-towards-
zero mode.
Parameters
i - [in] - short int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__host__ __device__ __nv_bfloat16 __short_as_bfloat16(const short int i)
Reinterprets bits in a signed short integer as a nv_bfloat16.
Reinterprets the bits in the signed short integer i as a nv_bfloat16 floating-point number.
5.5. Bfloat16 Precision Conversion and Data Movement
189
Parameters
i - [in] - short int. Is only being read.
Returns
nv_bfloat16
▶ The reinterpreted value.
__device__ void __stcg(__nv_bfloat16 *const ptr, const __nv_bfloat16 value)
Generates a st.global.cg store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stcg(__nv_bfloat162 *const ptr, const __nv_bfloat162 value)
Generates a st.global.cg store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stcs(__nv_bfloat16 *const ptr, const __nv_bfloat16 value)
Generates a st.global.cs store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stcs(__nv_bfloat162 *const ptr, const __nv_bfloat162 value)
Generates a st.global.cs store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stwb(__nv_bfloat16 *const ptr, const __nv_bfloat16 value)
Generates a st.global.wb store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stwb(__nv_bfloat162 *const ptr, const __nv_bfloat162 value)
Generates a st.global.wb store instruction.
Parameters
▶ ptr - [out] - memory location
190
Chapter 5. Bfloat16 Precision Intrinsics
▶ value - [in] - the value to be stored
__device__ void __stwt(__nv_bfloat162 *const ptr, const __nv_bfloat162 value)
Generates a st.global.wt store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ void __stwt(__nv_bfloat16 *const ptr, const __nv_bfloat16 value)
Generates a st.global.wt store instruction.
Parameters
▶ ptr - [out] - memory location
▶ value - [in] - the value to be stored
__device__ __nv_bfloat16 __uint2bfloat16_rd(const unsigned int i)
Convert an unsigned integer to a nv_bfloat16 in round-down mode.
Convert the unsigned integer value i to a nv_bfloat16 floating-point value in round-down mode.
Parameters
i - [in] - unsigned int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__host__ __device__ __nv_bfloat16 __uint2bfloat16_rn(const unsigned int i)
Convert an unsigned integer to a nv_bfloat16 in round-to-nearest-even mode.
Convert the unsigned integer value i to a nv_bfloat16 floating-point value in round-to-nearest-
even mode.
Parameters
i - [in] - unsigned int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __uint2bfloat16_ru(const unsigned int i)
Convert an unsigned integer to a nv_bfloat16 in round-up mode.
Convert the unsigned integer value i to a nv_bfloat16 floating-point value in round-up mode.
Parameters
i - [in] - unsigned int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
5.5. Bfloat16 Precision Conversion and Data Movement
191
__device__ __nv_bfloat16 __uint2bfloat16_rz(const unsigned int i)
Convert an unsigned integer to a nv_bfloat16 in round-towards-zero mode.
Convert the unsigned integer value i to a nv_bfloat16 floating-point value in round-towards-zero
mode.
Parameters
i - [in] - unsigned int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __ull2bfloat16_rd(const unsigned long long int i)
Convert an unsigned 64-bit integer to a nv_bfloat16 in round-down mode.
Convert the unsigned 64-bit integer value i to a nv_bfloat16 floating-point value in round-down
mode.
Parameters
i - [in] - unsigned long long int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__host__ __device__ __nv_bfloat16 __ull2bfloat16_rn(const unsigned long long int i)
Convert an unsigned 64-bit integer to a nv_bfloat16 in round-to-nearest-even mode.
Convert the unsigned 64-bit integer value i to a nv_bfloat16 floating-point value in round-to-
nearest-even mode.
Parameters
i - [in] - unsigned long long int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __ull2bfloat16_ru(const unsigned long long int i)
Convert an unsigned 64-bit integer to a nv_bfloat16 in round-up mode.
Convert the unsigned 64-bit integer value i to a nv_bfloat16 floating-point value in round-up
mode.
Parameters
i - [in] - unsigned long long int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __ull2bfloat16_rz(const unsigned long long int i)
Convert an unsigned 64-bit integer to a nv_bfloat16 in round-towards-zero mode.
Convert the unsigned 64-bit integer value i to a nv_bfloat16 floating-point value in round-
towards-zero mode.
Parameters
i - [in] - unsigned long long int. Is only being read.
192
Chapter 5. Bfloat16 Precision Intrinsics
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __ushort2bfloat16_rd(const unsigned short int i)
Convert an unsigned short integer to a nv_bfloat16 in round-down mode.
Convert the unsigned short integer value i to a nv_bfloat16 floating-point value in round-down
mode.
Parameters
i - [in] - unsigned short int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__host__ __device__ __nv_bfloat16 __ushort2bfloat16_rn(const unsigned short int i)
Convert an unsigned short integer to a nv_bfloat16 in round-to-nearest-even mode.
Convert the unsigned short integer value i to a nv_bfloat16 floating-point value in round-to-
nearest-even mode.
Parameters
i - [in] - unsigned short int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __ushort2bfloat16_ru(const unsigned short int i)
Convert an unsigned short integer to a nv_bfloat16 in round-up mode.
Convert the unsigned short integer value i to a nv_bfloat16 floating-point value in round-up
mode.
Parameters
i - [in] - unsigned short int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
__device__ __nv_bfloat16 __ushort2bfloat16_rz(const unsigned short int i)
Convert an unsigned short integer to a nv_bfloat16 in round-towards-zero mode.
Convert the unsigned short integer value i to a nv_bfloat16 floating-point value in round-
towards-zero mode.
Parameters
i - [in] - unsigned short int. Is only being read.
Returns
nv_bfloat16
▶ i converted to nv_bfloat16.
5.5. Bfloat16 Precision Conversion and Data Movement
193
__host__ __device__ __nv_bfloat16 __ushort_as_bfloat16(const unsigned short int i)
Reinterprets bits in an unsigned short integer as a nv_bfloat16.
Reinterprets the bits in the unsigned short integer i as a nv_bfloat16 floating-point number.
Parameters
i - [in] - unsigned short int. Is only being read.
Returns
nv_bfloat16
▶ The reinterpreted value.
__host__ __device__ __nv_bfloat162 make_bfloat162(const __nv_bfloat16 x, const __nv_bfloat16 y)
Vector function, combines two nv_bfloat16 numbers into one nv_bfloat162 number.
Combines two input nv_bfloat16 number x and y into one nv_bfloat162 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] - nv_bfloat16. Is only being read.
▶ y - [in] - nv_bfloat16. Is only being read.
Returns
__nv_bfloat162
▶ The __nv_bfloat162 vector with one half equal to x and the other to y.
5.6. Bfloat162 Arithmetic Functions
To use these functions, include the header file cuda_bf16.h in your program.
Functions
__host__ __device__ __nv_bfloat162 __h2div(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector division in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat162 __habs2(const __nv_bfloat162 a)
Calculates the absolute value of both halves of the input nv_bfloat162 number and returns the
result.
__host__ __device__ __nv_bfloat162 __hadd2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector addition in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat162 __hadd2_rn(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector addition in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat162 __hadd2_sat(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector addition in round-to-nearest-even mode, with saturation to [0.0,
1.0].
__device__ __nv_bfloat162 __hcmadd(const __nv_bfloat162 a, const __nv_bfloat162 b, const
__nv_bfloat162 c)
Performs fast complex multiply-accumulate.
194
Chapter 5. Bfloat16 Precision Intrinsics
__device__
__nv_bfloat162
__hfma2(const
__nv_bfloat162 a, const
__nv_bfloat162 b, const
__nv_bfloat162 c)
Performs nv_bfloat162 vector fused multiply-add in round-to-nearest-even mode.
__device__ __nv_bfloat162 __hfma2_relu(const __nv_bfloat162 a, const __nv_bfloat162 b, const
__nv_bfloat162 c)
Performs nv_bfloat162 vector fused multiply-add in round-to-nearest-even mode with relu
saturation.
__device__ __nv_bfloat162 __hfma2_sat(const __nv_bfloat162 a, const __nv_bfloat162 b, const
__nv_bfloat162 c)
Performs nv_bfloat162 vector fused multiply-add in round-to-nearest-even mode, with satu-
ration to [0.0, 1.0].
__host__ __device__ __nv_bfloat162 __hmul2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector multiplication in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat162 __hmul2_rn(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector multiplication in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat162 __hmul2_sat(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector multiplication in round-to-nearest-even mode, with saturation
to [0.0, 1.0].
__host__ __device__ __nv_bfloat162 __hneg2(const __nv_bfloat162 a)
Negates both halves of the input nv_bfloat162 number and returns the result.
__host__ __device__ __nv_bfloat162 __hsub2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector subtraction in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat162 __hsub2_rn(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector subtraction in round-to-nearest-even mode.
__host__ __device__ __nv_bfloat162 __hsub2_sat(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector subtraction in round-to-nearest-even mode, with saturation to
[0.0, 1.0].
__device__ __nv_bfloat162 atomicAdd(__nv_bfloat162 *const address, const __nv_bfloat162 val)
Vector add val to the value stored at address in global or shared memory, and writes this value
back to address .
__host__ __device__ __nv_bfloat162 operator*(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 multiplication operation.
__host__ __device__ __nv_bfloat162 & operator*=(__nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 compound assignment with multiplication operation.
__host__ __device__ __nv_bfloat162 operator+(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 addition operation.
__host__ __device__ __nv_bfloat162 operator+(const __nv_bfloat162 &h)
Implements packed nv_bfloat16 unary plus operator, returns input value.
__host__ __device__ __nv_bfloat162 operator++(__nv_bfloat162 &h, const int ignored)
Performs packed nv_bfloat16 postfix increment operation.
__host__ __device__ __nv_bfloat162 & operator++(__nv_bfloat162 &h)
Performs packed nv_bfloat16 prefix increment operation.
__host__ __device__ __nv_bfloat162 & operator+=(__nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 compound assignment with addition operation.
5.6. Bfloat162 Arithmetic Functions
195
__host__ __device__ __nv_bfloat162 operator-(const __nv_bfloat162 &h)
Implements packed nv_bfloat16 unary minus operator.
__host__ __device__ __nv_bfloat162 operator-(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 subtraction operation.
__host__ __device__ __nv_bfloat162 operator-(__nv_bfloat162 &h, const int ignored)
Performs packed nv_bfloat16 postfix decrement operation.
__host__ __device__ __nv_bfloat162 & operator-(__nv_bfloat162 &h)
Performs packed nv_bfloat16 prefix decrement operation.
__host__ __device__ __nv_bfloat162 & operator-=(__nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 compound assignment with subtraction operation.
__host__ __device__ __nv_bfloat162 operator/(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 division operation.
__host__ __device__ __nv_bfloat162 & operator/=(__nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 compound assignment with division operation.
5.6.1. Functions
__host__ __device__ __nv_bfloat162 __h2div(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector division in round-to-nearest-even mode.
Divides nv_bfloat162 input vector a by input vector b in round-to-nearest-even mode.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise division of a with b.
__host__ __device__ __nv_bfloat162 __habs2(const __nv_bfloat162 a)
Calculates the absolute value of both halves of the input nv_bfloat162 number and returns the
result.
Calculates the absolute value of both halves of the input nv_bfloat162 number and returns the
result.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
bfloat2
▶ Returns a with the absolute value of both halves.
__host__ __device__ __nv_bfloat162 __hadd2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector addition in round-to-nearest-even mode.
Performs nv_bfloat162 vector add of inputs a and b, in round-to-nearest-even mode.
196
Chapter 5. Bfloat16 Precision Intrinsics
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The sum of vectors a and b.
__host__ __device__ __nv_bfloat162 __hadd2_rn(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector addition in round-to-nearest-even mode.
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The sum of vectors a and b.
__host__ __device__ __nv_bfloat162 __hadd2_sat(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector addition in round-to-nearest-even mode, with saturation to [0.0,
1.0].
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The sum of a and b, with respect to saturation.
__device__ __nv_bfloat162 __hcmadd(const __nv_bfloat162 a, const __nv_bfloat162 b, const
__nv_bfloat162 c)
Performs fast complex multiply-accumulate.
Interprets vector nv_bfloat162 input pairs a, b, and c as complex numbers in nv_bfloat16
precision and performs complex multiply-accumulate operation: a*b + c
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
▶ c - [in] - nv_bfloat162. Is only being read.
5.6. Bfloat162 Arithmetic Functions
197
Returns
nv_bfloat162
▶ The result of complex multiply-accumulate operation on complex numbers a, b,
and c
__device__ __nv_bfloat162 __hfma2(const __nv_bfloat162 a, const __nv_bfloat162 b, const
__nv_bfloat162 c)
Performs nv_bfloat162 vector fused multiply-add in round-to-nearest-even mode.
Performs nv_bfloat162 vector multiply on inputs a and b, then performs a nv_bfloat162
vector add of the result with c, rounding the result once in round-to-nearest-even mode.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
▶ c - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The result of elementwise fused multiply-add operation on vectors a, b, and c.
__device__ __nv_bfloat162 __hfma2_relu(const __nv_bfloat162 a, const __nv_bfloat162 b, const
__nv_bfloat162 c)
Performs nv_bfloat162 vector fused multiply-add in round-to-nearest-even mode with relu
saturation.
Performs nv_bfloat162 vector multiply on inputs a and b, then performs a nv_bfloat162
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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
▶ c - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The result of elementwise fused multiply-add operation on vectors a, b, and c
with relu saturation.
__device__ __nv_bfloat162 __hfma2_sat(const __nv_bfloat162 a, const __nv_bfloat162 b, const
__nv_bfloat162 c)
Performs nv_bfloat162 vector fused multiply-add in round-to-nearest-even mode, with satu-
ration to [0.0, 1.0].
Performs nv_bfloat162 vector multiply on inputs a and b, then performs a nv_bfloat162
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
198
Chapter 5. Bfloat16 Precision Intrinsics
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
▶ c - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The result of elementwise fused multiply-add operation on vectors a, b, and c,
with respect to saturation.
__host__ __device__ __nv_bfloat162 __hmul2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector multiplication in round-to-nearest-even mode.
Performs nv_bfloat162 vector multiplication of inputs a and b, in round-to-nearest-even mode.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The result of elementwise multiplying the vectors a and b.
__host__ __device__ __nv_bfloat162 __hmul2_rn(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector multiplication in round-to-nearest-even mode.
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The result of elementwise multiplying the vectors a and b.
__host__ __device__ __nv_bfloat162 __hmul2_sat(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector multiplication in round-to-nearest-even mode, with saturation
to [0.0, 1.0].
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The result of elementwise multiplication of vectors a and b, with respect to
saturation.
5.6. Bfloat162 Arithmetic Functions
199
__host__ __device__ __nv_bfloat162 __hneg2(const __nv_bfloat162 a)
Negates both halves of the input nv_bfloat162 number and returns the result.
Negates both halves of the input nv_bfloat162 number a and returns the result.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ Returns a with both halves negated.
__host__ __device__ __nv_bfloat162 __hsub2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector subtraction in round-to-nearest-even mode.
Subtracts nv_bfloat162 input vector b from input vector a in round-to-nearest-even mode.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The subtraction of vector b from a.
__host__ __device__ __nv_bfloat162 __hsub2_rn(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector subtraction in round-to-nearest-even mode.
Subtracts nv_bfloat162 input vector b from input vector a in round-to-nearest-even mode.
Prevents floating-point contractions of mul+sub into fma.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The subtraction of vector b from a.
__host__ __device__ __nv_bfloat162 __hsub2_sat(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector subtraction in round-to-nearest-even mode, with saturation to
[0.0, 1.0].
Subtracts nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
200
Chapter 5. Bfloat16 Precision Intrinsics
▶ The subtraction of vector b from a, with respect to saturation.
__device__ __nv_bfloat162 atomicAdd(__nv_bfloat162 *const address, const __nv_bfloat162 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 nv_bfloat16
elements; the entire __nv_bfloat162 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 9.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] - __nv_bfloat162*. An address in global or shared memory.
▶ val - [in] - __nv_bfloat162. The value to be added.
Returns
__nv_bfloat162
▶ The old value read from address.
__host__ __device__ __nv_bfloat162 operator*(const __nv_bfloat162 &lh, const __nv_bfloat162
&rh)
Performs packed nv_bfloat16 multiplication operation.
See also __hmul2(__nv_bfloat162, __nv_bfloat162)
__host__ __device__ __nv_bfloat162 &operator*=(__nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 compound assignment with multiplication operation.
__host__ __device__ __nv_bfloat162 operator+(const __nv_bfloat162 &lh, const __nv_bfloat162
&rh)
Performs packed nv_bfloat16 addition operation.
See also __hadd2(__nv_bfloat162, __nv_bfloat162)
__host__ __device__ __nv_bfloat162 operator+(const __nv_bfloat162 &h)
Implements packed nv_bfloat16 unary plus operator, returns input value.
__host__ __device__ __nv_bfloat162 operator++(__nv_bfloat162 &h, const int ignored)
Performs packed nv_bfloat16 postfix increment operation.
__host__ __device__ __nv_bfloat162 &operator++(__nv_bfloat162 &h)
Performs packed nv_bfloat16 prefix increment operation.
__host__ __device__ __nv_bfloat162 &operator+=(__nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 compound assignment with addition operation.
__host__ __device__ __nv_bfloat162 operator-(const __nv_bfloat162 &h)
Implements packed nv_bfloat16 unary minus operator.
See also __hneg2(__nv_bfloat162)
5.6. Bfloat162 Arithmetic Functions
201
__host__ __device__ __nv_bfloat162 operator-(const __nv_bfloat162 &lh, const __nv_bfloat162
&rh)
Performs packed nv_bfloat16 subtraction operation.
See also __hsub2(__nv_bfloat162, __nv_bfloat162)
__host__ __device__ __nv_bfloat162 operator--(__nv_bfloat162 &h, const int ignored)
Performs packed nv_bfloat16 postfix decrement operation.
__host__ __device__ __nv_bfloat162 &operator--(__nv_bfloat162 &h)
Performs packed nv_bfloat16 prefix decrement operation.
__host__ __device__ __nv_bfloat162 &operator-=(__nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 compound assignment with subtraction operation.
__host__ __device__ __nv_bfloat162 operator∕(const __nv_bfloat162 &lh, const __nv_bfloat162
&rh)
Performs packed nv_bfloat16 division operation.
See also __h2div(__nv_bfloat162, __nv_bfloat162)
__host__ __device__ __nv_bfloat162 &operator∕=(__nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 compound assignment with division operation.
5.7. Bfloat162 Comparison Functions
To use these functions, include the header file cuda_bf16.h in your program.
Functions
__host__ __device__ bool __hbeq2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector if-equal comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ bool __hbequ2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered if-equal comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ bool __hbge2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-equal comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ bool __hbgeu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-equal comparison and returns boolean true
if both nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ bool __hbgt2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-than comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ bool __hbgtu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-than comparison and returns boolean true if
both nv_bfloat16 results are true, boolean false otherwise.
202
Chapter 5. Bfloat16 Precision Intrinsics
__host__ __device__ bool __hble2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-equal comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ bool __hbleu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-equal comparison and returns boolean true if
both nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ bool __hblt2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-than comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ bool __hbltu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-than comparison and returns boolean true if
both nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ bool __hbne2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector not-equal comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ bool __hbneu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered not-equal comparison and returns boolean true if
both nv_bfloat16 results are true, boolean false otherwise.
__host__ __device__ __nv_bfloat162 __heq2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector if-equal comparison.
__host__ __device__ unsigned int __heq2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector if-equal comparison.
__host__ __device__ __nv_bfloat162 __hequ2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered if-equal comparison.
__host__ __device__ unsigned int __hequ2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered if-equal comparison.
__host__ __device__ __nv_bfloat162 __hge2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-equal comparison.
__host__ __device__ unsigned int __hge2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-equal comparison.
__host__ __device__ __nv_bfloat162 __hgeu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-equal comparison.
__host__ __device__ unsigned int __hgeu2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-equal comparison.
__host__ __device__ __nv_bfloat162 __hgt2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-than comparison.
__host__ __device__ unsigned int __hgt2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-than comparison.
__host__ __device__ __nv_bfloat162 __hgtu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-than comparison.
__host__ __device__ unsigned int __hgtu2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-than comparison.
__host__ __device__ __nv_bfloat162 __hisnan2(const __nv_bfloat162 a)
Determine whether nv_bfloat162 argument is a NaN.
5.7. Bfloat162 Comparison Functions
203
__host__ __device__ __nv_bfloat162 __hle2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-equal comparison.
__host__ __device__ unsigned int __hle2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-equal comparison.
__host__ __device__ __nv_bfloat162 __hleu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-equal comparison.
__host__ __device__ unsigned int __hleu2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-equal comparison.
__host__ __device__ __nv_bfloat162 __hlt2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-than comparison.
__host__ __device__ unsigned int __hlt2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-than comparison.
__host__ __device__ __nv_bfloat162 __hltu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-than comparison.
__host__ __device__ unsigned int __hltu2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-than comparison.
__host__ __device__ __nv_bfloat162 __hmax2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Calculates nv_bfloat162 vector maximum of two inputs.
__host__ __device__ __nv_bfloat162 __hmax2_nan(const __nv_bfloat162 a, const __nv_bfloat162 b)
Calculates nv_bfloat162 vector maximum of two inputs, NaNs pass through.
__host__ __device__ __nv_bfloat162 __hmin2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Calculates nv_bfloat162 vector minimum of two inputs.
__host__ __device__ __nv_bfloat162 __hmin2_nan(const __nv_bfloat162 a, const __nv_bfloat162 b)
Calculates nv_bfloat162 vector minimum of two inputs, NaNs pass through.
__host__ __device__ __nv_bfloat162 __hne2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector not-equal comparison.
__host__ __device__ unsigned int __hne2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector not-equal comparison.
__host__ __device__ __nv_bfloat162 __hneu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered not-equal comparison.
__host__ __device__ unsigned int __hneu2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered not-equal comparison.
__host__ __device__ bool operator!=(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 unordered compare not-equal operation.
__host__ __device__ bool operator<(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 ordered less-than compare operation.
__host__ __device__ bool operator<=(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 ordered less-or-equal compare operation.
__host__ __device__ bool operator==(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 ordered compare equal operation.
__host__ __device__ bool operator>(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 ordered greater-than compare operation.
204
Chapter 5. Bfloat16 Precision Intrinsics
__host__ __device__ bool operator>=(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 ordered greater-or-equal compare operation.
5.7.1. Functions
__host__ __device__ bool __hbeq2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector if-equal comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector if-equal comparison of inputs a and b. The bool result is set to
true only if both nv_bfloat16 if-equal comparisons evaluate to true, or false otherwise. NaN
inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of if-equal comparison of vectors a and b are
true;
▶ false otherwise.
__host__ __device__ bool __hbequ2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered if-equal comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector if-equal comparison of inputs a and b. The bool result is set to
true only if both nv_bfloat16 if-equal comparisons evaluate to true, or false otherwise. NaN
inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of unordered if-equal comparison of vectors
a and b are true;
▶ false otherwise.
__host__ __device__ bool __hbge2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-equal comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector greater-equal comparison of inputs a and b. The bool result
is set to true only if both nv_bfloat16 greater-equal comparisons evaluate to true, or false
otherwise. NaN inputs generate false results.
5.7. Bfloat162 Comparison Functions
205
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of greater-equal comparison of vectors a and
b are true;
▶ false otherwise.
__host__ __device__ bool __hbgeu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-equal comparison and returns boolean true
if both nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector greater-equal comparison of inputs a and b. The bool result
is set to true only if both nv_bfloat16 greater-equal comparisons evaluate to true, or false
otherwise. NaN inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of unordered greater-equal comparison of
vectors a and b are true;
▶ false otherwise.
__host__ __device__ bool __hbgt2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-than comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector greater-than comparison of inputs a and b. The bool result is set
to true only if both nv_bfloat16 greater-than comparisons evaluate to true, or false otherwise.
NaN inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of greater-than comparison of vectors a and
b are true;
▶ false otherwise.
__host__ __device__ bool __hbgtu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-than comparison and returns boolean true if
both nv_bfloat16 results are true, boolean false otherwise.
206
Chapter 5. Bfloat16 Precision Intrinsics
Performs nv_bfloat162 vector greater-than comparison of inputs a and b. The bool result is set
to true only if both nv_bfloat16 greater-than comparisons evaluate to true, or false otherwise.
NaN inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of unordered greater-than comparison of vec-
tors a and b are true;
▶ false otherwise.
__host__ __device__ bool __hble2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-equal comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector less-equal comparison of inputs a and b. The bool result is set
to true only if both nv_bfloat16 less-equal comparisons evaluate to true, or false otherwise.
NaN inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of less-equal comparison of vectors a and b
are true;
▶ false otherwise.
__host__ __device__ bool __hbleu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-equal comparison and returns boolean true if
both nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector less-equal comparison of inputs a and b. The bool result is set
to true only if both nv_bfloat16 less-equal comparisons evaluate to true, or false otherwise.
NaN inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of unordered less-equal comparison of vec-
tors a and b are true;
▶ false otherwise.
5.7. Bfloat162 Comparison Functions
207
__host__ __device__ bool __hblt2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-than comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector less-than comparison of inputs a and b. The bool result is set to
true only if both nv_bfloat16 less-than comparisons evaluate to true, or false otherwise. NaN
inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of less-than comparison of vectors a and b
are true;
▶ false otherwise.
__host__ __device__ bool __hbltu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-than comparison and returns boolean true if
both nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector less-than comparison of inputs a and b. The bool result is set to
true only if both nv_bfloat16 less-than comparisons evaluate to true, or false otherwise. NaN
inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of unordered less-than comparison of vectors
a and b are true;
▶ false otherwise.
__host__ __device__ bool __hbne2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector not-equal comparison and returns boolean true if both
nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector not-equal comparison of inputs a and b. The bool result is set to
true only if both nv_bfloat16 not-equal comparisons evaluate to true, or false otherwise. NaN
inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
208
Chapter 5. Bfloat16 Precision Intrinsics
▶ true if both nv_bfloat16 results of not-equal comparison of vectors a and b
are true,
▶ false otherwise.
__host__ __device__ bool __hbneu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered not-equal comparison and returns boolean true if
both nv_bfloat16 results are true, boolean false otherwise.
Performs nv_bfloat162 vector not-equal comparison of inputs a and b. The bool result is set to
true only if both nv_bfloat16 not-equal comparisons evaluate to true, or false otherwise. NaN
inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
bool
▶ true if both nv_bfloat16 results of unordered not-equal comparison of vectors
a and b are true;
▶ false otherwise.
__host__ __device__ __nv_bfloat162 __heq2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector if-equal comparison.
Performs nv_bfloat162 vector if-equal comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector result of if-equal comparison of vectors a and b.
__host__ __device__ unsigned int __heq2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector if-equal comparison.
Performs nv_bfloat162 vector if-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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of if-equal comparison of vectors a and b.
5.7. Bfloat162 Comparison Functions
209
__host__ __device__ __nv_bfloat162 __hequ2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered if-equal comparison.
Performs nv_bfloat162 vector if-equal comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector result of unordered if-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hequ2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered if-equal comparison.
Performs nv_bfloat162 vector if-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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered if-equal comparison of vectors a and b.
__host__ __device__ __nv_bfloat162 __hge2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-equal comparison.
Performs nv_bfloat162 vector greater-equal comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector result of greater-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hge2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-equal comparison.
Performs nv_bfloat162 vector greater-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] - nv_bfloat162. Is only being read.
210
Chapter 5. Bfloat16 Precision Intrinsics
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of greater-equal comparison of vectors a and b.
__host__ __device__ __nv_bfloat162 __hgeu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-equal comparison.
Performs nv_bfloat162 vector greater-equal comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The nv_bfloat162 vector result of unordered greater-equal comparison of
vectors a and b.
__host__ __device__ unsigned int __hgeu2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-equal comparison.
Performs nv_bfloat162 vector greater-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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered greater-equal comparison of vectors a and
b.
__host__ __device__ __nv_bfloat162 __hgt2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-than comparison.
Performs nv_bfloat162 vector greater-than comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector result of greater-than comparison of vectors a and b.
5.7. Bfloat162 Comparison Functions
211
__host__ __device__ unsigned int __hgt2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector greater-than comparison.
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of greater-than comparison of vectors a and b.
__host__ __device__ __nv_bfloat162 __hgtu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-than comparison.
Performs nv_bfloat162 vector greater-than comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The nv_bfloat162 vector result of unordered greater-than comparison of vec-
tors a and b.
__host__ __device__ unsigned int __hgtu2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered greater-than comparison.
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered greater-than comparison of vectors a and
b.
__host__ __device__ __nv_bfloat162 __hisnan2(const __nv_bfloat162 a)
Determine whether nv_bfloat162 argument is a NaN.
Determine whether each nv_bfloat16 of input nv_bfloat162 number a is a NaN.
Parameters
a - [in] - nv_bfloat162. Is only being read.
212
Chapter 5. Bfloat16 Precision Intrinsics
Returns
nv_bfloat162
▶ The nv_bfloat162 with the corresponding nv_bfloat16 results set to 1.0 for
NaN, 0.0 otherwise.
__host__ __device__ __nv_bfloat162 __hle2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-equal comparison.
Performs nv_bfloat162 vector less-equal comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The nv_bfloat162 result of less-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hle2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-equal comparison.
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of less-equal comparison of vectors a and b.
__host__ __device__ __nv_bfloat162 __hleu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-equal comparison.
Performs nv_bfloat162 vector less-equal comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector result of unordered less-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hleu2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-equal comparison.
Performs nv_bfloat162 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.
5.7. Bfloat162 Comparison Functions
213
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered less-equal comparison of vectors a and b.
__host__ __device__ __nv_bfloat162 __hlt2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-than comparison.
Performs nv_bfloat162 vector less-than comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The nv_bfloat162 vector result of less-than comparison of vectors a and b.
__host__ __device__ unsigned int __hlt2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector less-than comparison.
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of less-than comparison of vectors a and b.
__host__ __device__ __nv_bfloat162 __hltu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-than comparison.
Performs nv_bfloat162 vector less-than comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector result of unordered less-than comparison of vectors a and b.
214
Chapter 5. Bfloat16 Precision Intrinsics
__host__ __device__ unsigned int __hltu2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered less-than comparison.
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered less-than comparison of vectors a and b.
__host__ __device__ __nv_bfloat162 __hmax2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Calculates nv_bfloat162 vector maximum of two inputs.
Calculates nv_bfloat162 vector max(a, b). Elementwise nv_bfloat16 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
__host__ __device__ __nv_bfloat162 __hmax2_nan(const __nv_bfloat162 a, const __nv_bfloat162 b)
Calculates nv_bfloat162 vector maximum of two inputs, NaNs pass through.
Calculates nv_bfloat162 vector max(a, b). Elementwise nv_bfloat16 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 maximum of vectors a and b, with NaNs pass through
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
5.7. Bfloat162 Comparison Functions
215
__host__ __device__ __nv_bfloat162 __hmin2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Calculates nv_bfloat162 vector minimum of two inputs.
Calculates nv_bfloat162 vector min(a, b). Elementwise nv_bfloat16 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
__host__ __device__ __nv_bfloat162 __hmin2_nan(const __nv_bfloat162 a, const __nv_bfloat162 b)
Calculates nv_bfloat162 vector minimum of two inputs, NaNs pass through.
Calculates nv_bfloat162 vector min(a, b). Elementwise nv_bfloat16 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
__host__ __device__ __nv_bfloat162 __hne2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector not-equal comparison.
Performs nv_bfloat162 vector not-equal comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector result of not-equal comparison of vectors a and b.
216
Chapter 5. Bfloat16 Precision Intrinsics
__host__ __device__ unsigned int __hne2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector not-equal comparison.
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of not-equal comparison of vectors a and b.
__host__ __device__ __nv_bfloat162 __hneu2(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered not-equal comparison.
Performs nv_bfloat162 vector not-equal comparison of inputs a and b. The corresponding
nv_bfloat16 results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
Parameters
▶ a - [in] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector result of unordered not-equal comparison of vectors a and b.
__host__ __device__ unsigned int __hneu2_mask(const __nv_bfloat162 a, const __nv_bfloat162 b)
Performs nv_bfloat162 vector unordered not-equal comparison.
Performs nv_bfloat162 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] - nv_bfloat162. Is only being read.
▶ b - [in] - nv_bfloat162. Is only being read.
Returns
unsigned int
▶ The vector mask result of unordered not-equal comparison of vectors a and b.
__host__ __device__ bool operator!=(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 unordered compare not-equal operation.
See also __hbneu2(__nv_bfloat162, __nv_bfloat162)
__host__ __device__ bool operator<(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 ordered less-than compare operation.
See also __hblt2(__nv_bfloat162, __nv_bfloat162)
5.7. Bfloat162 Comparison Functions
217
__host__ __device__ bool operator<=(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 ordered less-or-equal compare operation.
See also __hble2(__nv_bfloat162, __nv_bfloat162)
__host__ __device__ bool operator==(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 ordered compare equal operation.
See also __hbeq2(__nv_bfloat162, __nv_bfloat162)
__host__ __device__ bool operator>(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 ordered greater-than compare operation.
See also __hbgt2(__nv_bfloat162, __nv_bfloat162)
__host__ __device__ bool operator>=(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh)
Performs packed nv_bfloat16 ordered greater-or-equal compare operation.
See also __hbge2(__nv_bfloat162, __nv_bfloat162)
5.8. Bfloat162 Math Functions
To use these functions, include the header file cuda_bf16.h in your program.
Functions
__device__ __nv_bfloat162 h2ceil(const __nv_bfloat162 h)
Calculate nv_bfloat162 vector ceiling of the input argument.
__device__ __nv_bfloat162 h2cos(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector cosine in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2exp(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector exponential function in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2exp10(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector decimal exponential function in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2exp2(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector binary exponential function in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2floor(const __nv_bfloat162 h)
Calculate the largest integer less than or equal to h .
__device__ __nv_bfloat162 h2log(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector natural logarithm in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2log10(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector decimal logarithm in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2log2(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector binary logarithm in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2rcp(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector reciprocal in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2rint(const __nv_bfloat162 h)
Round input to nearest integer value in nv_bfloat16 floating-point number.
218
Chapter 5. Bfloat16 Precision Intrinsics
__device__ __nv_bfloat162 h2rsqrt(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector reciprocal square root in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2sin(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector sine in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2sqrt(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector square root in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2tanh(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector hyperbolic tangent function in round-to-nearest-even mode.
__device__ __nv_bfloat162 h2tanh_approx(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector approximate hyperbolic tangent function.
__device__ __nv_bfloat162 h2trunc(const __nv_bfloat162 h)
Truncate nv_bfloat162 vector input argument to the integral part.
5.8.1. Functions
__device__ __nv_bfloat162 h2ceil(const __nv_bfloat162 h)
Calculate nv_bfloat162 vector ceiling of the input argument.
For each component of vector h compute the smallest integer value not less than h.
Parameters
h - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector of smallest integers not less than h.
__device__ __nv_bfloat162 h2cos(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector cosine in round-to-nearest-even mode.
Calculates nv_bfloat162 cosine of input vector 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_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise cosine on vector a.
__device__ __nv_bfloat162 h2exp(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector exponential function in round-to-nearest-even mode.
Calculates nv_bfloat162 exponential function of input vector a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
5.8. Bfloat162 Math Functions
219
▶ The elementwise exponential function on vector a.
__device__ __nv_bfloat162 h2exp10(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector decimal exponential function in round-to-nearest-even mode.
Calculates nv_bfloat162 decimal exponential function of input vector a in round-to-nearest-
even mode.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise decimal exponential function on vector a.
__device__ __nv_bfloat162 h2exp2(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector binary exponential function in round-to-nearest-even mode.
Calculates nv_bfloat162 binary exponential function of input vector a in round-to-nearest-
even mode.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise binary exponential function on vector a.
__device__ __nv_bfloat162 h2floor(const __nv_bfloat162 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.
Parameters
h - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector of largest integers which is less than or equal to h.
__device__ __nv_bfloat162 h2log(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector natural logarithm in round-to-nearest-even mode.
Calculates nv_bfloat162 natural logarithm of input vector a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise natural logarithm on vector a.
__device__ __nv_bfloat162 h2log10(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector decimal logarithm in round-to-nearest-even mode.
Calculates nv_bfloat162 decimal logarithm of input vector a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat162. Is only being read.
220
Chapter 5. Bfloat16 Precision Intrinsics
Returns
nv_bfloat162
▶ The elementwise decimal logarithm on vector a.
__device__ __nv_bfloat162 h2log2(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector binary logarithm in round-to-nearest-even mode.
Calculates nv_bfloat162 binary logarithm of input vector a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise binary logarithm on vector a.
__device__ __nv_bfloat162 h2rcp(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector reciprocal in round-to-nearest-even mode.
Calculates nv_bfloat162 reciprocal of input vector a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise reciprocal on vector a.
__device__ __nv_bfloat162 h2rint(const __nv_bfloat162 h)
Round input to nearest integer value in nv_bfloat16 floating-point number.
Round each component of nv_bfloat162 vector h to the nearest integer value in nv_bfloat16
floating-point format, with bfloat16way cases rounded to the nearest even integer value.
Parameters
h - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The vector of rounded integer values.
__device__ __nv_bfloat162 h2rsqrt(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector reciprocal square root in round-to-nearest-even mode.
Calculates nv_bfloat162 reciprocal square root of input vector a in round-to-nearest-even
mode.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise reciprocal square root on vector a.
__device__ __nv_bfloat162 h2sin(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector sine in round-to-nearest-even mode.
Calculates nv_bfloat162 sine of input vector a in round-to-nearest-even mode.
5.8. Bfloat162 Math Functions
221
NOTE: this function’s implementation calls sinf(float) function and is exposed to compiler opti-
mizations. Specifically, --use_fast_math flag changes sinf(float) into an intrinsic __sinf(float),
which has less accurate numeric behavior.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise sine on vector a.
__device__ __nv_bfloat162 h2sqrt(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector square root in round-to-nearest-even mode.
Calculates nv_bfloat162 square root of input vector a in round-to-nearest-even mode.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise square root on vector a.
__device__ __nv_bfloat162 h2tanh(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector hyperbolic tangent function in round-to-nearest-even mode.
Calculates nv_bfloat162 hyperbolic tangent function of input vector a in round-to-nearest-
even mode.
See also:
htanh(__nv_bfloat16) for further details.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise hyperbolic tangent function on vector a.
__device__ __nv_bfloat162 h2tanh_approx(const __nv_bfloat162 a)
Calculates nv_bfloat162 vector approximate hyperbolic tangent function.
Calculates nv_bfloat162 approximate hyperbolic tangent function of input vector a. This op-
eration uses HW acceleration on devices of compute capability 9.x and higher.
See also:
htanh_approx(__nv_bfloat16) for further details.
Parameters
a - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The elementwise approximate hyperbolic tangent function on vector a.
222
Chapter 5. Bfloat16 Precision Intrinsics
__device__ __nv_bfloat162 h2trunc(const __nv_bfloat162 h)
Truncate nv_bfloat162 vector input argument to the integral part.
Round each component of vector h to the nearest integer value that does not exceed h in mag-
nitude.
Parameters
h - [in] - nv_bfloat162. Is only being read.
Returns
nv_bfloat162
▶ The truncated h.
Groups
Bfloat16 Arithmetic Constants
To use these constants, include the header file cuda_bf16.h in your program.
Bfloat16 Arithmetic Functions
To use these functions, include the header file cuda_bf16.h in your program.
Bfloat16 Comparison Functions
To use these functions, include the header file cuda_bf16.h in your program.
Bfloat16 Math Functions
To use these functions, include the header file cuda_bf16.h in your program.
Bfloat16 Precision Conversion and Data Movement
To use these functions, include the header file cuda_bf16.h in your program.
Bfloat162 Arithmetic Functions
To use these functions, include the header file cuda_bf16.h in your program.
Bfloat162 Comparison Functions
To use these functions, include the header file cuda_bf16.h in your program.
Bfloat162 Math Functions
To use these functions, include the header file cuda_bf16.h in your program.
Structs
__nv_bfloat16
nv_bfloat16 datatype
__nv_bfloat162
nv_bfloat162 datatype
__nv_bfloat162_raw
__nv_bfloat162_raw data type
__nv_bfloat16_raw
__nv_bfloat16_raw data type
5.8. Bfloat162 Math Functions
223
Typedefs
nv_bfloat16
This datatype is meant to be the first-class or fundamental implementation of the bfloat16 num-
bers format.
nv_bfloat162
This datatype is meant to be the first-class or fundamental implementation of type for pairs of
bfloat16 numbers.
5.9. Typedefs
typedef __nv_bfloat16 nv_bfloat16
This datatype is meant to be the first-class or fundamental implementation of the bfloat16 num-
bers 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 __nv_bfloat162 nv_bfloat162
This datatype is meant to be the first-class or fundamental implementation of type for pairs of
bfloat16 numbers.
Should be implemented in the compiler in the future. Current implementation is a simple typedef
to a respective user-level type with underscores.
224
Chapter 5. Bfloat16 Precision Intrinsics
Chapter 6. Single Precision
Mathematical Functions
This section describes single precision mathematical functions.
To use these functions, you do not need to include any additional header file in your program.
Functions
__device__ float acosf(float x)
Calculate the arc cosine of the input argument.
__device__ float acoshf(float x)
Calculate the nonnegative inverse hyperbolic cosine of the input argument.
__device__ float asinf(float x)
Calculate the arc sine of the input argument.
__device__ float asinhf(float x)
Calculate the inverse hyperbolic sine of the input argument.
__device__ float atan2f(float y, float x)
Calculate the arc tangent of the ratio of first and second input arguments.
__device__ float atanf(float x)
Calculate the arc tangent of the input argument.
__device__ float atanhf(float x)
Calculate the inverse hyperbolic tangent of the input argument.
__device__ float cbrtf(float x)
Calculate the cube root of the input argument.
__device__ float ceilf(float x)
Calculate ceiling of the input argument.
__device__ float copysignf(float x, float y)
Create value with given magnitude, copying sign of second value.
__device__ float cosf(float x)
Calculate the cosine of the input argument.
__device__ float coshf(float x)
Calculate the hyperbolic cosine of the input argument.
__device__ float cospif(float x)
Calculate the cosine of the input argument ×π .
225
__device__ float cyl_bessel_i0f(float x)
Calculate the value of the regular modified cylindrical Bessel function of order 0 for the input
argument.
__device__ float cyl_bessel_i1f(float x)
Calculate the value of the regular modified cylindrical Bessel function of order 1 for the input
argument.
__device__ float erfcf(float x)
Calculate the complementary error function of the input argument.
__device__ float erfcinvf(float x)
Calculate the inverse complementary error function of the input argument.
__device__ float erfcxf(float x)
Calculate the scaled complementary error function of the input argument.
__device__ float erff(float x)
Calculate the error function of the input argument.
__device__ float erfinvf(float x)
Calculate the inverse error function of the input argument.
__device__ float exp10f(float x)
Calculate the base 10 exponential of the input argument.
__device__ float exp2f(float x)
Calculate the base 2 exponential of the input argument.
__device__ float expf(float x)
Calculate the base e exponential of the input argument.
__device__ float expm1f(float x)
Calculate the base e exponential of the input argument, minus 1.
__device__ float fabsf(float x)
Calculate the absolute value of its argument.
__device__ float fdimf(float x, float y)
Compute the positive difference between x and y .
__device__ float fdividef(float x, float y)
Divide two floating-point values.
__device__ float floorf(float x)
Calculate the largest integer less than or equal to x .
__device__ float fmaf(float x, float y, float z)
Compute x × y + z as a single operation.
__device__ float fmaxf(float x, float y)
Determine the maximum numeric value of the arguments.
__device__ float fminf(float x, float y)
Determine the minimum numeric value of the arguments.
__device__ float fmodf(float x, float y)
Calculate the floating-point remainder of x / y .
__device__ float frexpf(float x, int *nptr)
Extract mantissa and exponent of a floating-point value.
__device__ float hypotf(float x, float y)
Calculate the square root of the sum of squares of two arguments.
226
Chapter 6. Single Precision Mathematical Functions
__device__ int ilogbf(float x)
Compute the unbiased integer exponent of the argument.
__device__ __RETURN_TYPE isfinite(float a)
Determine whether argument is finite.
__device__ __RETURN_TYPE isinf(float a)
Determine whether argument is infinite.
__device__ __RETURN_TYPE isnan(float a)
Determine whether argument is a NaN.
__device__ float j0f(float x)
Calculate the value of the Bessel function of the first kind of order 0 for the input argument.
__device__ float j1f(float x)
Calculate the value of the Bessel function of the first kind of order 1 for the input argument.
__device__ float jnf(int n, float x)
Calculate the value of the Bessel function of the first kind of order n for the input argument.
__device__ float ldexpf(float x, int exp)
Calculate the value of x · 2exp .
__device__ float lgammaf(float x)
Calculate the natural logarithm of the absolute value of the gamma function of the input argu-
ment.
__device__ long long int llrintf(float x)
Round input to nearest integer value.
__device__ long long int llroundf(float x)
Round to nearest integer value.
__device__ float log10f(float x)
Calculate the base 10 logarithm of the input argument.
__device__ float log1pf(float x)
Calculate the value of loge(1 + x) .
__device__ float log2f(float x)
Calculate the base 2 logarithm of the input argument.
__device__ float logbf(float x)
Calculate the floating-point representation of the exponent of the input argument.
__device__ float logf(float x)
Calculate the natural logarithm of the input argument.
__device__ long int lrintf(float x)
Round input to nearest integer value.
__device__ long int lroundf(float x)
Round to nearest integer value.
__device__ float max(const float a, const float b)
Calculate the maximum value of the input float arguments.
__device__ float min(const float a, const float b)
Calculate the minimum value of the input float arguments.
__device__ float modff(float x, float *iptr)
Break down the input argument into fractional and integral parts.
227
__device__ float nanf(const char *tagp)
Returns "Not a Number" value.
__device__ float nearbyintf(float x)
Round the input argument to the nearest integer.
__device__ float nextafterf(float x, float y)
Return next representable single-precision floating-point value after argument x in the direction
of y .
__device__ float norm3df(float a, float b, float c)
Calculate the square root of the sum of squares of three coordinates of the argument.
__device__ float norm4df(float a, float b, float c, float d)
Calculate the square root of the sum of squares of four coordinates of the argument.
__device__ float normcdff(float x)
Calculate the standard normal cumulative distribution function.
__device__ float normcdfinvf(float x)
Calculate the inverse of the standard normal cumulative distribution function.
__device__ float normf(int dim, float const *p)
Calculate the square root of the sum of squares of any number of coordinates.
__device__ float powf(float x, float y)
Calculate the value of first argument to the power of second argument.
__device__ float rcbrtf(float x)
Calculate reciprocal cube root function.
__device__ float remainderf(float x, float y)
Compute single-precision floating-point remainder.
__device__ float remquof(float x, float y, int *quo)
Compute single-precision floating-point remainder and part of quotient.
__device__ float rhypotf(float x, float y)
Calculate one over the square root of the sum of squares of two arguments.
__device__ float rintf(float x)
Round input to nearest integer value in floating-point.
__device__ float rnorm3df(float a, float b, float c)
Calculate one over the square root of the sum of squares of three coordinates.
__device__ float rnorm4df(float a, float b, float c, float d)
Calculate one over the square root of the sum of squares of four coordinates.
__device__ float rnormf(int dim, float const *p)
Calculate the reciprocal of square root of the sum of squares of any number of coordinates.
__device__ float roundf(float x)
Round to nearest integer value in floating-point.
__device__ float rsqrtf(float x)
Calculate the reciprocal of the square root of the input argument.
__device__ float scalblnf(float x, long int n)
Scale floating-point input by integer power of two.
__device__ float scalbnf(float x, int n)
Scale floating-point input by integer power of two.
228
Chapter 6. Single Precision Mathematical Functions
__device__ __RETURN_TYPE signbit(float a)
Return the sign bit of the input.
__device__ void sincosf(float x, float *sptr, float *cptr)
Calculate the sine and cosine of the first input argument.
__device__ void sincospif(float x, float *sptr, float *cptr)
Calculate the sine and cosine of the first input argument ×π .
__device__ float sinf(float x)
Calculate the sine of the input argument.
__device__ float sinhf(float x)
Calculate the hyperbolic sine of the input argument.
__device__ float sinpif(float x)
Calculate the sine of the input argument ×π .
__device__ float sqrtf(float x)
Calculate the square root of the input argument.
__device__ float tanf(float x)
Calculate the tangent of the input argument.
__device__ float tanhf(float x)
Calculate the hyperbolic tangent of the input argument.
__device__ float tgammaf(float x)
Calculate the gamma function of the input argument.
__device__ float truncf(float x)
Truncate input argument to the integral part.
__device__ float y0f(float x)
Calculate the value of the Bessel function of the second kind of order 0 for the input argument.
__device__ float y1f(float x)
Calculate the value of the Bessel function of the second kind of order 1 for the input argument.
__device__ float ynf(int n, float x)
Calculate the value of the Bessel function of the second kind of order n for the input argument.
6.1. Functions
__device__ float acosf(float x)
Calculate the arc cosine of the input argument.
Calculate the principal value of the arc cosine of the input argument x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
Result will be in radians, in the interval [0, π ] for x inside [-1, +1].
▶ acosf(1) returns +0.
6.1. Functions
229
▶ acosf(x) returns NaN for x outside [-1, +1].
▶ acosf(NaN) returns NaN.
__device__ float acoshf(float x)
Calculate the nonnegative inverse hyperbolic cosine of the input argument.
Calculate the nonnegative inverse hyperbolic cosine of the input argument x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
Result will be in the interval [0, +∞ ].
▶ acoshf(1) returns 0.
▶ acoshf(x) returns NaN for x in the interval [ -∞ , 1).
▶ acoshf( +∞ ) returns +∞.
▶ acoshf(NaN) returns NaN.
__device__ float asinf(float x)
Calculate the arc sine of the input argument.
Calculate the principal value of the arc sine of the input argument x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
Result will be in radians, in the interval [- π/2 , + π/2 ] for x inside [-1, +1].
▶ asinf( ±0 ) returns ±0.
▶ asinf(x) returns NaN for x outside [-1, +1].
▶ asinf(NaN) returns NaN.
__device__ float asinhf(float x)
Calculate the inverse hyperbolic sine of the input argument.
Calculate the inverse hyperbolic sine of the input argument x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
▶ asinhf( ±0 ) returns ±0.
▶ asinhf( ±∞ ) returns ±∞.
▶ asinhf(NaN) returns NaN.
230
Chapter 6. Single Precision Mathematical Functions
__device__ float atan2f(float y, float x)
Calculate the arc tangent of the ratio of first and second input arguments.
Calculate the principal value of the arc tangent of the ratio of first and second input arguments
y / x. The quadrant of the result is determined by the signs of inputs y and x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
Result will be in radians, in the interval [- π , + π ].
▶ atan2f( ±0 , -0) returns ±π.
▶ atan2f( ±0 , +0) returns ±0.
▶ atan2f( ±0 , x) returns ±π for x < 0.
▶ atan2f( ±0 , x) returns ±0 for x > 0.
▶ atan2f(y, ±0 ) returns -π /2 for y < 0.
▶ atan2f(y, ±0 ) returns π /2 for y > 0.
▶ atan2f( ±y , -∞ ) returns ±π for finite y > 0.
▶ atan2f( ±y , +∞ ) returns ±0 for finite y > 0.
▶ atan2f( ±∞ , x) returns ±π /2 for finite x.
▶ atan2f( ±∞ , -∞ ) returns ±3π /4.
▶ atan2f( ±∞ , +∞ ) returns ±π /4.
▶ If either argument is NaN, NaN is returned.
__device__ float atanf(float x)
Calculate the arc tangent of the input argument.
Calculate the principal value of the arc tangent of the input argument x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
Result will be in radians, in the interval [- π/2 , + π/2 ].
▶ atanf( ±0 ) returns ±0.
▶ atanf( ±∞ ) returns ±π /2.
▶ atanf(NaN) returns NaN.
__device__ float atanhf(float x)
Calculate the inverse hyperbolic tangent of the input argument.
Calculate the inverse hyperbolic tangent of the input argument x.
6.1. Functions
231
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
▶ atanhf( ±0 ) returns ±0.
▶ atanhf( ±1 ) returns ±∞.
▶ atanhf(x) returns NaN for x outside interval [-1, 1].
▶ atanhf(NaN) returns NaN.
__device__ float cbrtf(float x)
Calculate the cube root of the input argument.
Calculate the cube root of x, x1/3.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
Returns x1/3.
▶ cbrtf( ±0 ) returns ±0.
▶ cbrtf( ±∞ ) returns ±∞.
▶ cbrtf(NaN) returns NaN.
__device__ float ceilf(float x)
Calculate ceiling of the input argument.
Compute the smallest integer value not less than x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
Returns ⌈x⌉ expressed as a floating-point number.
▶ ceilf( ±0 ) returns ±0.
▶ ceilf( ±∞ ) returns ±∞.
▶ ceilf(NaN) returns NaN.
__device__ float copysignf(float x, float y)
Create value with given magnitude, copying sign of second value.
Create a floating-point value with the magnitude x and the sign of y.
Returns
232
Chapter 6. Single Precision Mathematical Functions
▶ a value with the magnitude of x and the sign of y.
▶ copysignf(NaN, y) returns a NaN with the sign of y.
__device__ float cosf(float x)
Calculate the cosine of the input argument.
Calculate the cosine of the input argument x (measured in radians).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Note: This function is affected by the use_fast_math compiler flag. See the CUDA C++ Pro-
gramming Guide, Mathematical Functions Appendix, Intrinsic Functions section for a complete
list of functions affected.
Returns
▶ cosf( ±0 ) returns 1.
▶ cosf( ±∞ ) returns NaN.
▶ cosf(NaN) returns NaN.
__device__ float coshf(float x)
Calculate the hyperbolic cosine of the input argument.
Calculate the hyperbolic cosine of the input argument x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
▶ coshf( ±0 ) returns 1.
▶ coshf( ±∞ ) returns +∞.
▶ coshf(NaN) returns NaN.
__device__ float cospif(float x)
Calculate the cosine of the input argument ×π.
Calculate the cosine of x ×π (measured in radians), where x is the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
6.1. Functions
233
▶ cospif( ±0 ) returns 1.
▶ cospif( ±∞ ) returns NaN.
▶ cospif(NaN) returns NaN.
__device__ float cyl_bessel_i0f(float x)
Calculate the value of the regular modified cylindrical Bessel function of order 0 for the input
argument.
Calculate the value of the regular modified cylindrical Bessel function of order 0 for the input
argument x, I0(x).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
Returns the value of the regular modified cylindrical Bessel function of order 0.
▶ cyl_bessel_i0f( ±0) returns +1.
▶ cyl_bessel_i0f( ±∞) returns +∞.
▶ cyl_bessel_i0f(NaN) returns NaN.
__device__ float cyl_bessel_i1f(float x)
Calculate the value of the regular modified cylindrical Bessel function of order 1 for the input
argument.
Calculate the value of the regular modified cylindrical Bessel function of order 1 for the input
argument x, I1(x).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
Returns the value of the regular modified cylindrical Bessel function of order 1.
▶ cyl_bessel_i1f( ±0) returns ±0.
▶ cyl_bessel_i1f( ±∞) returns ±∞.
▶ cyl_bessel_i1f(NaN) returns NaN.
__device__ float erfcf(float x)
Calculate the complementary error function of the input argument.
Calculate the complementary error function of the input argument x, 1 - erf(x).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Single-Precision Floating-Point Functions section.
Returns
234
Chapter 6. Single Precision Mathematical Functions
|
||
|
|
|