|
|
__device__ unsigned int __vcmpltu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: a < b ? 0xff : 0.
__device__ unsigned int __vcmpne2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed comparison: a != b ? 0xffff : 0.
__device__ unsigned int __vcmpne4(unsigned int a, unsigned int b)
Performs per-byte (un)signed comparison: a != b ? 0xff : 0.
__device__ unsigned int __vhaddu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned average computation.
__device__ unsigned int __vhaddu4(unsigned int a, unsigned int b)
Computes per-byte unsigned average.
__host__ __device__ unsigned int __viaddmax_s16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword max(a + b, c)
__host__ __device__ unsigned int __viaddmax_s16x2_relu(const unsigned int a, const unsigned int
b, const unsigned int c)
Performs per-halfword max(max(a + b, c), 0)
__host__ __device__ int __viaddmax_s32(const int a, const int b, const int c)
Computes max(a + b, c)
__host__ __device__ int __viaddmax_s32_relu(const int a, const int b, const int c)
Computes max(max(a + b, c), 0)
__host__ __device__ unsigned int __viaddmax_u16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword max(a + b, c)
__host__ __device__ unsigned int __viaddmax_u32(const unsigned int a, const unsigned int b, const
unsigned int c)
Computes max(a + b, c)
__host__ __device__ unsigned int __viaddmin_s16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword min(a + b, c)
__host__ __device__ unsigned int __viaddmin_s16x2_relu(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword max(min(a + b, c), 0)
__host__ __device__ int __viaddmin_s32(const int a, const int b, const int c)
Computes min(a + b, c)
__host__ __device__ int __viaddmin_s32_relu(const int a, const int b, const int c)
Computes max(min(a + b, c), 0)
__host__ __device__ unsigned int __viaddmin_u16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword min(a + b, c)
__host__ __device__ unsigned int __viaddmin_u32(const unsigned int a, const unsigned int b, const
unsigned int c)
Computes min(a + b, c)
__host__ __device__ unsigned int __vibmax_s16x2(const unsigned int a, const unsigned int b, bool
*const pred_hi, bool *const pred_lo)
395
Performs per-halfword max(a, b), also sets the value pointed to by pred_hi and pred_lo to the
per-halfword result of (a >= b).
__host__ __device__ int __vibmax_s32(const int a, const int b, bool *const pred)
Computes max(a, b), also sets the value pointed to by pred to (a >= b).
__host__ __device__ unsigned int __vibmax_u16x2(const unsigned int a, const unsigned int b, bool
*const pred_hi, bool *const pred_lo)
Performs per-halfword max(a, b), also sets the value pointed to by pred_hi and pred_lo to the
per-halfword result of (a >= b).
__host__ __device__ unsigned int __vibmax_u32(const unsigned int a, const unsigned int b, bool
*const pred)
Computes max(a, b), also sets the value pointed to by pred to (a >= b).
__host__ __device__ unsigned int __vibmin_s16x2(const unsigned int a, const unsigned int b, bool
*const pred_hi, bool *const pred_lo)
Performs per-halfword min(a, b), also sets the value pointed to by pred_hi and pred_lo to the
per-halfword result of (a <= b).
__host__ __device__ int __vibmin_s32(const int a, const int b, bool *const pred)
Computes min(a, b), also sets the value pointed to by pred to (a <= b).
__host__ __device__ unsigned int __vibmin_u16x2(const unsigned int a, const unsigned int b, bool
*const pred_hi, bool *const pred_lo)
Performs per-halfword min(a, b), also sets the value pointed to by pred_hi and pred_lo to the
per-halfword result of (a <= b).
__host__ __device__ unsigned int __vibmin_u32(const unsigned int a, const unsigned int b, bool
*const pred)
Computes min(a, b), also sets the value pointed to by pred to (a <= b).
__host__ __device__ unsigned int __vimax3_s16x2(const unsigned int a, const unsigned int b, const
unsigned int c)
Performs per-halfword max(max(a, b), c)
__host__ __device__ unsigned int __vimax3_s16x2_relu(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword max(max(max(a, b), c), 0)
__host__ __device__ int __vimax3_s32(const int a, const int b, const int c)
Computes max(max(a, b), c)
__host__ __device__ int __vimax3_s32_relu(const int a, const int b, const int c)
Computes max(max(max(a, b), c), 0)
__host__ __device__ unsigned int __vimax3_u16x2(const unsigned int a, const unsigned int b, const
unsigned int c)
Performs per-halfword max(max(a, b), c)
__host__ __device__ unsigned int __vimax3_u32(const unsigned int a, const unsigned int b, const
unsigned int c)
Computes max(max(a, b), c)
__host__ __device__ unsigned int __vimax_s16x2_relu(const unsigned int a, const unsigned int b)
Performs per-halfword max(max(a, b), 0)
__host__ __device__ int __vimax_s32_relu(const int a, const int b)
Computes max(max(a, b), 0)
396
Chapter 14. SIMD Intrinsics
__host__ __device__ unsigned int __vimin3_s16x2(const unsigned int a, const unsigned int b, const
unsigned int c)
Performs per-halfword min(min(a, b), c)
__host__ __device__ unsigned int __vimin3_s16x2_relu(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword max(min(min(a, b), c), 0)
__host__ __device__ int __vimin3_s32(const int a, const int b, const int c)
Computes min(min(a, b), c)
__host__ __device__ int __vimin3_s32_relu(const int a, const int b, const int c)
Computes max(min(min(a, b), c), 0)
__host__ __device__ unsigned int __vimin3_u16x2(const unsigned int a, const unsigned int b, const
unsigned int c)
Performs per-halfword min(min(a, b), c)
__host__ __device__ unsigned int __vimin3_u32(const unsigned int a, const unsigned int b, const
unsigned int c)
Computes min(min(a, b), c)
__host__ __device__ unsigned int __vimin_s16x2_relu(const unsigned int a, const unsigned int b)
Performs per-halfword max(min(a, b), 0)
__host__ __device__ int __vimin_s32_relu(const int a, const int b)
Computes max(min(a, b), 0)
__device__ unsigned int __vmaxs2(unsigned int a, unsigned int b)
Performs per-halfword signed maximum computation.
__device__ unsigned int __vmaxs4(unsigned int a, unsigned int b)
Computes per-byte signed maximum.
__device__ unsigned int __vmaxu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned maximum computation.
__device__ unsigned int __vmaxu4(unsigned int a, unsigned int b)
Computes per-byte unsigned maximum.
__device__ unsigned int __vmins2(unsigned int a, unsigned int b)
Performs per-halfword signed minimum computation.
__device__ unsigned int __vmins4(unsigned int a, unsigned int b)
Computes per-byte signed minimum.
__device__ unsigned int __vminu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned minimum computation.
__device__ unsigned int __vminu4(unsigned int a, unsigned int b)
Computes per-byte unsigned minimum.
__device__ unsigned int __vneg2(unsigned int a)
Computes per-halfword negation.
__device__ unsigned int __vneg4(unsigned int a)
Performs per-byte negation.
__device__ unsigned int __vnegss2(unsigned int a)
Computes per-halfword negation with signed saturation.
__device__ unsigned int __vnegss4(unsigned int a)
Performs per-byte negation with signed saturation.
397
__device__ unsigned int __vsads2(unsigned int a, unsigned int b)
Performs per-halfword sum of absolute difference of signed.
__device__ unsigned int __vsads4(unsigned int a, unsigned int b)
Computes per-byte sum of abs difference of signed.
__device__ unsigned int __vsadu2(unsigned int a, unsigned int b)
Computes per-halfword sum of abs diff of unsigned.
__device__ unsigned int __vsadu4(unsigned int a, unsigned int b)
Computes per-byte sum of abs difference of unsigned.
__device__ unsigned int __vseteq2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed comparison: returns 1 if both parts compare equal.
__device__ unsigned int __vseteq4(unsigned int a, unsigned int b)
Performs per-byte (un)signed comparison: returns 1 if all 4 pairs compare equal.
__device__ unsigned int __vsetges2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: returns 1 if both parts compare greater than or equal.
__device__ unsigned int __vsetges4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: returns 1 if all 4 pairs compare greater than or equal.
__device__ unsigned int __vsetgeu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: returns 1 if both parts compare greater than or
equal.
__device__ unsigned int __vsetgeu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: returns 1 if all 4 pairs compare greater than or equal.
__device__ unsigned int __vsetgts2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: returns 1 if both parts compare greater than.
__device__ unsigned int __vsetgts4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: returns 1 if all 4 pairs compare greater than.
__device__ unsigned int __vsetgtu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: returns 1 if both parts compare greater than.
__device__ unsigned int __vsetgtu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: returns 1 if all 4 pairs compare greater than.
__device__ unsigned int __vsetles2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: returns 1 if both parts compare less than or equal.
__device__ unsigned int __vsetles4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: returns 1 if all 4 pairs compare less than or equal.
__device__ unsigned int __vsetleu2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: returns 1 if both parts compare less than or equal.
__device__ unsigned int __vsetleu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: returns 1 if all 4 pairs compare less than or equal.
__device__ unsigned int __vsetlts2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: returns 1 if both parts compare less than.
__device__ unsigned int __vsetlts4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: returns 1 if all 4 pairs compare less than.
__device__ unsigned int __vsetltu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: returns 1 if both parts compare less than.
398
Chapter 14. SIMD Intrinsics
__device__ unsigned int __vsetltu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: returns 1 if all 4 pairs compare less than.
__device__ unsigned int __vsetne2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed comparison: returns 1 if both parts compare not equal.
__device__ unsigned int __vsetne4(unsigned int a, unsigned int b)
Performs per-byte (un)signed comparison: returns 1 if all 4 pairs compare not equal.
__device__ unsigned int __vsub2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed subtraction, with wrap-around: a - b.
__device__ unsigned int __vsub4(unsigned int a, unsigned int b)
Performs per-byte subtraction: a - b.
__device__ unsigned int __vsubss2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed subtraction, with signed saturation: a - b.
__device__ unsigned int __vsubss4(unsigned int a, unsigned int b)
Performs per-byte subtraction with signed saturation: a - b.
__device__ unsigned int __vsubus2(unsigned int a, unsigned int b)
Performs per-halfword subtraction with unsigned saturation: a - b.
__device__ unsigned int __vsubus4(unsigned int a, unsigned int b)
Performs per-byte subtraction with unsigned saturation: a - b.
14.1. Functions
__device__ unsigned int __vabs2(unsigned int a)
Computes per-halfword absolute value: |a|.
Splits 4 bytes of argument into 2 parts, each consisting of 2 bytes, then computes absolute value
for each of parts. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vabs4(unsigned int a)
Computes per-byte absolute value: |a|.
Splits argument by bytes. Computes absolute value of each byte. Partial results are recombined
and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vabsdiffs2(unsigned int a, unsigned int b)
Computes per-halfword absolute difference of signed integer: |a - b|.
Splits 4 bytes of each into 2 parts, each consisting of 2 bytes. For corresponding parts function
computes absolute difference. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
14.1. Functions
399
__device__ unsigned int __vabsdiffs4(unsigned int a, unsigned int b)
Computes per-byte absolute difference of signed integer: |a - b|.
Splits 4 bytes of each into 4 parts, each consisting of 1 byte. For corresponding parts function
computes absolute difference. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vabsdiffu2(unsigned int a, unsigned int b)
Computes per-halfword absolute difference of unsigned integer: |a - b|.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function computes absolute difference. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vabsdiffu4(unsigned int a, unsigned int b)
Computes per-byte absolute difference of unsigned integer: |a - b|.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function computes absolute difference. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vabsss2(unsigned int a)
Computes per-halfword absolute value with signed saturation: |a|.
Splits 4 bytes of argument into 2 parts, each consisting of 2 bytes, then computes absolute value
with signed saturation for each of parts. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vabsss4(unsigned int a)
Computes per-byte absolute value with signed saturation: |a|.
Splits 4 bytes of argument into 4 parts, each consisting of 1 byte, then computes absolute value
with signed saturation for each of parts. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vadd2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed addition, with wrap-around: a + b.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes, then performs unsigned
addition on corresponding parts. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vadd4(unsigned int a, unsigned int b)
Performs per-byte (un)signed addition: a + b.
400
Chapter 14. SIMD Intrinsics
Splits ‘a’ into 4 bytes, then performs unsigned addition on each of these bytes with the corre-
sponding byte from ‘b’, ignoring overflow. Partial results are recombined and returned as un-
signed int.
Returns
Returns computed value.
__device__ unsigned int __vaddss2(unsigned int a, unsigned int b)
Performs per-halfword addition with signed saturation: a + b.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes, then performs addition
with signed saturation on corresponding parts. Partial results are recombined and returned as
unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vaddss4(unsigned int a, unsigned int b)
Performs per-byte addition with signed saturation: a + b.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte, then performs addition
with signed saturation on corresponding parts. Partial results are recombined and returned as
unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vaddus2(unsigned int a, unsigned int b)
Performs per-halfword addition with unsigned saturation: a + b.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes, then performs addition
with unsigned saturation on corresponding parts.
Returns
Returns computed value.
__device__ unsigned int __vaddus4(unsigned int a, unsigned int b)
Performs per-byte addition with unsigned saturation: a + b.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte, then performs addition
with unsigned saturation on corresponding parts.
Returns
Returns computed value.
__device__ unsigned int __vavgs2(unsigned int a, unsigned int b)
Performs per-halfword signed rounded average computation.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes, then computes signed
rounded average of corresponding parts. Partial results are recombined and returned as un-
signed int.
Returns
Returns computed value.
__device__ unsigned int __vavgs4(unsigned int a, unsigned int b)
Computes per-byte signed rounded average.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. then computes signed
rounded average of corresponding parts. Partial results are recombined and returned as un-
signed int.
14.1. Functions
401
Returns
Returns computed value.
__device__ unsigned int __vavgu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned rounded average computation.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes, then computes un-
signed rounded average of corresponding parts. Partial results are recombined and returned as
unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vavgu4(unsigned int a, unsigned int b)
Performs per-byte unsigned rounded average.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. then computes un-
signed rounded average of corresponding parts. Partial results are recombined and returned as
unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vcmpeq2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed comparison: a == b ? 0xffff : 0.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts result is ffff if they are equal, and 0000 otherwise. For example __vcmpeq2(0x1234aba5,
0x1234aba6) returns 0xffff0000.
Returns
Returns 0xffff computed value.
__device__ unsigned int __vcmpeq4(unsigned int a, unsigned int b)
Performs per-byte (un)signed comparison: a == b ? 0xff : 0.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
result is ff if they are equal, and 00 otherwise. For example __vcmpeq4(0x1234aba5, 0x1234aba6)
returns 0xffffff00.
Returns
Returns 0xff if a = b, else returns 0.
__device__ unsigned int __vcmpges2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: a >= b ? 0xffff : 0.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
result is ffff if ‘a’ part >= ‘b’ part, and 0000 otherwise. For example __vcmpges2(0x1234aba5,
0x1234aba6) returns 0xffff0000.
Returns
Returns 0xffff if a >= b, else returns 0.
__device__ unsigned int __vcmpges4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: a >= b ? 0xff : 0.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding
parts result is ff if ‘a’ part >= ‘b’ part, and 00 otherwise. For example __vcmpges4(0x1234aba5,
0x1234aba6) returns 0xffffff00.
Returns
Returns 0xff if a >= b, else returns 0.
402
Chapter 14. SIMD Intrinsics
__device__ unsigned int __vcmpgeu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: a >= b ? 0xffff : 0.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
result is ffff if ‘a’ part >= ‘b’ part, and 0000 otherwise. For example __vcmpgeu2(0x1234aba5,
0x1234aba6) returns 0xffff0000.
Returns
Returns 0xffff if a >= b, else returns 0.
__device__ unsigned int __vcmpgeu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: a >= b ? 0xff : 0.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding
parts result is ff if ‘a’ part >= ‘b’ part, and 00 otherwise. For example __vcmpgeu4(0x1234aba5,
0x1234aba6) returns 0xffffff00.
Returns
Returns 0xff if a >= b, else returns 0.
__device__ unsigned int __vcmpgts2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: a > b ? 0xffff : 0.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts result is ffff if ‘a’ part > ‘b’ part, and 0000 otherwise. For example __vcmpgts2(0x1234aba5,
0x1234aba6) returns 0x00000000.
Returns
Returns 0xffff if a > b, else returns 0.
__device__ unsigned int __vcmpgts4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: a > b ? 0xff : 0.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding
parts result is ff if ‘a’ part > ‘b’ part, and 00 otherwise. For example __vcmpgts4(0x1234aba5,
0x1234aba6) returns 0x00000000.
Returns
Returns 0xff if a > b, else returns 0.
__device__ unsigned int __vcmpgtu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: a > b ? 0xffff : 0.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts result is ffff if ‘a’ part > ‘b’ part, and 0000 otherwise. For example __vcmpgtu2(0x1234aba5,
0x1234aba6) returns 0x00000000.
Returns
Returns 0xffff if a > b, else returns 0.
__device__ unsigned int __vcmpgtu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: a > b ? 0xff : 0.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding
parts result is ff if ‘a’ part > ‘b’ part, and 00 otherwise. For example __vcmpgtu4(0x1234aba5,
0x1234aba6) returns 0x00000000.
Returns
Returns 0xff if a > b, else returns 0.
14.1. Functions
403
__device__ unsigned int __vcmples2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: a <= b ? 0xffff : 0.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
result is ffff if ‘a’ part <= ‘b’ part, and 0000 otherwise. For example __vcmples2(0x1234aba5,
0x1234aba6) returns 0xffffffff.
Returns
Returns 0xffff if a <= b, else returns 0.
__device__ unsigned int __vcmples4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: a <= b ? 0xff : 0.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding
parts result is ff if ‘a’ part <= ‘b’ part, and 00 otherwise. For example __vcmples4(0x1234aba5,
0x1234aba6) returns 0xffffffff.
Returns
Returns 0xff if a <= b, else returns 0.
__device__ unsigned int __vcmpleu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: a <= b ? 0xffff : 0.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
result is ffff if ‘a’ part <= ‘b’ part, and 0000 otherwise. For example __vcmpleu2(0x1234aba5,
0x1234aba6) returns 0xffffffff.
Returns
Returns 0xffff if a <= b, else returns 0.
__device__ unsigned int __vcmpleu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: a <= b ? 0xff : 0.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding
parts result is ff if ‘a’ part <= ‘b’ part, and 00 otherwise. For example __vcmpleu4(0x1234aba5,
0x1234aba6) returns 0xffffffff.
Returns
Returns 0xff if a <= b, else returns 0.
__device__ unsigned int __vcmplts2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: a < b ? 0xffff : 0.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts result is ffff if ‘a’ part < ‘b’ part, and 0000 otherwise. For example __vcmplts2(0x1234aba5,
0x1234aba6) returns 0x0000ffff.
Returns
Returns 0xffff if a < b, else returns 0.
__device__ unsigned int __vcmplts4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: a < b ? 0xff : 0.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding
parts result is ff if ‘a’ part < ‘b’ part, and 00 otherwise. For example __vcmplts4(0x1234aba5,
0x1234aba6) returns 0x000000ff.
Returns
Returns 0xff if a < b, else returns 0.
404
Chapter 14. SIMD Intrinsics
__device__ unsigned int __vcmpltu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: a < b ? 0xffff : 0.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts result is ffff if ‘a’ part < ‘b’ part, and 0000 otherwise. For example __vcmpltu2(0x1234aba5,
0x1234aba6) returns 0x0000ffff.
Returns
Returns 0xffff if a < b, else returns 0.
__device__ unsigned int __vcmpltu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: a < b ? 0xff : 0.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding
parts result is ff if ‘a’ part < ‘b’ part, and 00 otherwise. For example __vcmpltu4(0x1234aba5,
0x1234aba6) returns 0x000000ff.
Returns
Returns 0xff if a < b, else returns 0.
__device__ unsigned int __vcmpne2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed comparison: a != b ? 0xffff : 0.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts result is ffff if ‘a’ part != ‘b’ part, and 0000 otherwise. For example __vcmplts2(0x1234aba5,
0x1234aba6) returns 0x0000ffff.
Returns
Returns 0xffff if a != b, else returns 0.
__device__ unsigned int __vcmpne4(unsigned int a, unsigned int b)
Performs per-byte (un)signed comparison: a != b ? 0xff : 0.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding
parts result is ff if ‘a’ part != ‘b’ part, and 00 otherwise. For example __vcmplts4(0x1234aba5,
0x1234aba6) returns 0x000000ff.
Returns
Returns 0xff if a != b, else returns 0.
__device__ unsigned int __vhaddu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned average computation.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes, then computes unsigned
average of corresponding parts. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vhaddu4(unsigned int a, unsigned int b)
Computes per-byte unsigned average.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. then computes unsigned
average of corresponding parts. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__host__ __device__ unsigned int __viaddmax_s16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
14.1. Functions
405
Performs per-halfword max(a + b, c)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts are
interpreted as signed shorts. For corresponding parts function performs an add and compare:
max(a_part + b_part), c_part) Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__host__ __device__ unsigned int __viaddmax_s16x2_relu(const unsigned int a, const unsigned
int b, const unsigned int c)
Performs per-halfword max(max(a + b, c), 0)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts are
interpreted as signed shorts. For corresponding parts function performs an add, followed by a
max with relu: max(max(a_part + b_part), c_part), 0) Partial results are recombined and returned
as unsigned int.
Returns
Returns computed value.
__host__ __device__ int __viaddmax_s32(const int a, const int b, const int c)
Computes max(a + b, c)
Calculates the sum of signed integers a and b and takes the max with c.
Returns
Returns computed value.
__host__ __device__ int __viaddmax_s32_relu(const int a, const int b, const int c)
Computes max(max(a + b, c), 0)
Calculates the sum of signed integers a and b and takes the max with c. If the result is less than
0 then 0 is returned.
Returns
Returns computed value.
__host__ __device__ unsigned int __viaddmax_u16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword max(a + b, c)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts are
interpreted as unsigned shorts. For corresponding parts function performs an add and compare:
max(a_part + b_part), c_part) Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__host__ __device__ unsigned int __viaddmax_u32(const unsigned int a, const unsigned int b,
const unsigned int c)
Computes max(a + b, c)
Calculates the sum of unsigned integers a and b and takes the max with c.
Returns
Returns computed value.
406
Chapter 14. SIMD Intrinsics
__host__ __device__ unsigned int __viaddmin_s16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword min(a + b, c)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts are
interpreted as signed shorts. For corresponding parts function performs an add and compare:
min(a_part + b_part), c_part) Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__host__ __device__ unsigned int __viaddmin_s16x2_relu(const unsigned int a, const unsigned
int b, const unsigned int c)
Performs per-halfword max(min(a + b, c), 0)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts are
interpreted as signed shorts. For corresponding parts function performs an add, followed by a
min with relu: max(min(a_part + b_part), c_part), 0) Partial results are recombined and returned
as unsigned int.
Returns
Returns computed value.
__host__ __device__ int __viaddmin_s32(const int a, const int b, const int c)
Computes min(a + b, c)
Calculates the sum of signed integers a and b and takes the min with c.
Returns
Returns computed value.
__host__ __device__ int __viaddmin_s32_relu(const int a, const int b, const int c)
Computes max(min(a + b, c), 0)
Calculates the sum of signed integers a and b and takes the min with c. If the result is less than
0 then 0 is returned.
Returns
Returns computed value.
__host__ __device__ unsigned int __viaddmin_u16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword min(a + b, c)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts are
interpreted as unsigned shorts. For corresponding parts function performs an add and compare:
min(a_part + b_part), c_part) Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__host__ __device__ unsigned int __viaddmin_u32(const unsigned int a, const unsigned int b,
const unsigned int c)
Computes min(a + b, c)
Calculates the sum of unsigned integers a and b and takes the min with c.
Returns
Returns computed value.
14.1. Functions
407
__host__ __device__ unsigned int __vibmax_s16x2(const unsigned int a, const unsigned int b, bool
*const pred_hi, bool *const pred_lo)
Performs per-halfword max(a, b), also sets the value pointed to by pred_hi and pred_lo to the
per-halfword result of (a >= b).
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as signed shorts. For corresponding parts function performs a maximum ( =
max(a_part, b_part) ). Partial results are recombined and returned as unsigned int. Sets the
value pointed to by pred_hi to the value (a_high_part >= b_high_part). Sets the value pointed
to by pred_lo to the value (a_low_part >= b_low_part).
Returns
Returns computed values.
__host__ __device__ int __vibmax_s32(const int a, const int b, bool *const pred)
Computes max(a, b), also sets the value pointed to by pred to (a >= b).
Calculates the maximum of a and b of two signed ints. Also sets the value pointed to by pred to
the value (a >= b).
Returns
Returns computed values.
__host__ __device__ unsigned int __vibmax_u16x2(const unsigned int a, const unsigned int b, bool
*const pred_hi, bool *const pred_lo)
Performs per-halfword max(a, b), also sets the value pointed to by pred_hi and pred_lo to the
per-halfword result of (a >= b).
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as unsigned shorts. For corresponding parts function performs a maximum ( =
max(a_part, b_part) ). Partial results are recombined and returned as unsigned int. Sets the value
pointed to by pred_hi to the value (a_high_part >= b_high_part). Sets the value pointed to by
pred_lo to the value (a_low_part >= b_low_part).
Returns
Returns computed values.
__host__ __device__ unsigned int __vibmax_u32(const unsigned int a, const unsigned int b, bool
*const pred)
Computes max(a, b), also sets the value pointed to by pred to (a >= b).
Calculates the maximum of a and b of two unsigned ints. Also sets the value pointed to by pred
to the value (a >= b).
Returns
Returns computed values.
__host__ __device__ unsigned int __vibmin_s16x2(const unsigned int a, const unsigned int b, bool
*const pred_hi, bool *const pred_lo)
Performs per-halfword min(a, b), also sets the value pointed to by pred_hi and pred_lo to the
per-halfword result of (a <= b).
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as signed shorts. For corresponding parts function performs a maximum ( =
max(a_part, b_part) ). Partial results are recombined and returned as unsigned int. Sets the
value pointed to by pred_hi to the value (a_high_part <= b_high_part). Sets the value pointed
to by pred_lo to the value (a_low_part <= b_low_part).
Returns
Returns computed values.
408
Chapter 14. SIMD Intrinsics
__host__ __device__ int __vibmin_s32(const int a, const int b, bool *const pred)
Computes min(a, b), also sets the value pointed to by pred to (a <= b).
Calculates the minimum of a and b of two signed ints. Also sets the value pointed to by pred to
the value (a <= b).
Returns
Returns computed values.
__host__ __device__ unsigned int __vibmin_u16x2(const unsigned int a, const unsigned int b, bool
*const pred_hi, bool *const pred_lo)
Performs per-halfword min(a, b), also sets the value pointed to by pred_hi and pred_lo to the
per-halfword result of (a <= b).
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as unsigned shorts. For corresponding parts function performs a maximum ( =
max(a_part, b_part) ). Partial results are recombined and returned as unsigned int. Sets the value
pointed to by pred_hi to the value (a_high_part <= b_high_part). Sets the value pointed to by
pred_lo to the value (a_low_part <= b_low_part).
Returns
Returns computed values.
__host__ __device__ unsigned int __vibmin_u32(const unsigned int a, const unsigned int b, bool
*const pred)
Computes min(a, b), also sets the value pointed to by pred to (a <= b).
Calculates the minimum of a and b of two unsigned ints. Also sets the value pointed to by pred
to the value (a <= b).
Returns
Returns computed values.
__host__ __device__ unsigned int __vimax3_s16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword max(max(a, b), c)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as signed shorts. For corresponding parts function performs a 3-way max ( =
max(max(a_part, b_part), c_part) ). Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__host__ __device__ unsigned int __vimax3_s16x2_relu(const unsigned int a, const unsigned int
b, const unsigned int c)
Performs per-halfword max(max(max(a, b), c), 0)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts are
interpreted as signed shorts. For corresponding parts function performs a three-way max with
relu ( = max(a_part, b_part, c_part, 0) ). Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__host__ __device__ int __vimax3_s32(const int a, const int b, const int c)
Computes max(max(a, b), c)
Calculates the 3-way max of signed integers a, b and c.
14.1. Functions
409
Returns
Returns computed value.
__host__ __device__ int __vimax3_s32_relu(const int a, const int b, const int c)
Computes max(max(max(a, b), c), 0)
Calculates the maximum of three signed ints, if this is less than 0 then 0 is returned.
Returns
Returns computed value.
__host__ __device__ unsigned int __vimax3_u16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword max(max(a, b), c)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as unsigned shorts. For corresponding parts function performs a 3-way max ( =
max(max(a_part, b_part), c_part) ). Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__host__ __device__ unsigned int __vimax3_u32(const unsigned int a, const unsigned int b, const
unsigned int c)
Computes max(max(a, b), c)
Calculates the 3-way max of unsigned integers a, b and c.
Returns
Returns computed value.
__host__ __device__ unsigned int __vimax_s16x2_relu(const unsigned int a, const unsigned int
b)
Performs per-halfword max(max(a, b), 0)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as signed shorts. For corresponding parts function performs a max with relu ( =
max(a_part, b_part, 0) ). Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__host__ __device__ int __vimax_s32_relu(const int a, const int b)
Computes max(max(a, b), 0)
Calculates the maximum of a and b of two signed ints, if this is less than 0 then 0 is returned.
Returns
Returns computed value.
__host__ __device__ unsigned int __vimin3_s16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword min(min(a, b), c)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as signed shorts. For corresponding parts function performs a 3-way min ( =
min(min(a_part, b_part), c_part) ). Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
410
Chapter 14. SIMD Intrinsics
__host__ __device__ unsigned int __vimin3_s16x2_relu(const unsigned int a, const unsigned int
b, const unsigned int c)
Performs per-halfword max(min(min(a, b), c), 0)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as signed shorts. For corresponding parts function performs a three-way min
with relu ( = max(min(a_part, b_part, c_part), 0) ). Partial results are recombined and returned as
unsigned int.
Returns
Returns computed value.
__host__ __device__ int __vimin3_s32(const int a, const int b, const int c)
Computes min(min(a, b), c)
Calculates the 3-way min of signed integers a, b and c.
Returns
Returns computed value.
__host__ __device__ int __vimin3_s32_relu(const int a, const int b, const int c)
Computes max(min(min(a, b), c), 0)
Calculates the minimum of three signed ints, if this is less than 0 then 0 is returned.
Returns
Returns computed value.
__host__ __device__ unsigned int __vimin3_u16x2(const unsigned int a, const unsigned int b,
const unsigned int c)
Performs per-halfword min(min(a, b), c)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as unsigned shorts. For corresponding parts function performs a 3-way min ( =
min(min(a_part, b_part), c_part) ). Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__host__ __device__ unsigned int __vimin3_u32(const unsigned int a, const unsigned int b, const
unsigned int c)
Computes min(min(a, b), c)
Calculates the 3-way min of unsigned integers a, b and c.
Returns
Returns computed value.
__host__ __device__ unsigned int __vimin_s16x2_relu(const unsigned int a, const unsigned int
b)
Performs per-halfword max(min(a, b), 0)
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. These 2 byte parts
are interpreted as signed shorts. For corresponding parts function performs a min with relu ( =
max(min(a_part, b_part), 0) ). Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
14.1. Functions
411
__host__ __device__ int __vimin_s32_relu(const int a, const int b)
Computes max(min(a, b), 0)
Calculates the minimum of a and b of two signed ints, if this is less than 0 then 0 is returned.
Returns
Returns computed value.
__device__ unsigned int __vmaxs2(unsigned int a, unsigned int b)
Performs per-halfword signed maximum computation.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function computes signed maximum. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vmaxs4(unsigned int a, unsigned int b)
Computes per-byte signed maximum.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function computes signed maximum. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vmaxu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned maximum computation.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function computes unsigned maximum. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vmaxu4(unsigned int a, unsigned int b)
Computes per-byte unsigned maximum.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function computes unsigned maximum. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vmins2(unsigned int a, unsigned int b)
Performs per-halfword signed minimum computation.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function computes signed minimum. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vmins4(unsigned int a, unsigned int b)
Computes per-byte signed minimum.
412
Chapter 14. SIMD Intrinsics
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function computes signed minimum. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vminu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned minimum computation.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function computes unsigned minimum. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vminu4(unsigned int a, unsigned int b)
Computes per-byte unsigned minimum.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function computes unsigned minimum. Partial results are recombined and returned as unsigned
int.
Returns
Returns computed value.
__device__ unsigned int __vneg2(unsigned int a)
Computes per-halfword negation.
Splits 4 bytes of argument into 2 parts, each consisting of 2 bytes. For each part function com-
putes negation. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vneg4(unsigned int a)
Performs per-byte negation.
Splits 4 bytes of argument into 4 parts, each consisting of 1 byte. For each part function com-
putes negation. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vnegss2(unsigned int a)
Computes per-halfword negation with signed saturation.
Splits 4 bytes of argument into 2 parts, each consisting of 2 bytes. For each part function com-
putes negation. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vnegss4(unsigned int a)
Performs per-byte negation with signed saturation.
Splits 4 bytes of argument into 4 parts, each consisting of 1 byte. For each part function com-
putes negation. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
14.1. Functions
413
__device__ unsigned int __vsads2(unsigned int a, unsigned int b)
Performs per-halfword sum of absolute difference of signed.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts function computes absolute difference and sum it up. Partial results are recombined and
returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vsads4(unsigned int a, unsigned int b)
Computes per-byte sum of abs difference of signed.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function computes absolute difference and sum it up. Partial results are recombined and re-
turned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vsadu2(unsigned int a, unsigned int b)
Computes per-halfword sum of abs diff of unsigned.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function computes absolute differences and returns sum of those differences.
Returns
Returns computed value.
__device__ unsigned int __vsadu4(unsigned int a, unsigned int b)
Computes per-byte sum of abs difference of unsigned.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function computes absolute differences and returns sum of those differences.
Returns
Returns computed value.
__device__ unsigned int __vseteq2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed comparison: returns 1 if both parts compare equal.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function performs comparison ‘a’ part == ‘b’ part. If both equalities are satisfied, function returns
1.
Returns
Returns 1 if a = b, else returns 0.
__device__ unsigned int __vseteq4(unsigned int a, unsigned int b)
Performs per-byte (un)signed comparison: returns 1 if all 4 pairs compare equal.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs comparison ‘a’ part == ‘b’ part. If both equalities are satisfied, function returns
1.
Returns
Returns 1 if a = b, else returns 0.
__device__ unsigned int __vsetges2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: returns 1 if both parts compare greater than or equal.
414
Chapter 14. SIMD Intrinsics
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts function performs comparison ‘a’ part >= ‘b’ part. If both inequalities are satisfied, function
returns 1.
Returns
Returns 1 if a >= b, else returns 0.
__device__ unsigned int __vsetges4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: returns 1 if all 4 pairs compare greater than or equal.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs comparison ‘a’ part >= ‘b’ part. If both inequalities are satisfied, function re-
turns 1.
Returns
Returns 1 if a >= b, else returns 0.
__device__ unsigned int __vsetgeu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: returns 1 if both parts compare greater than or
equal.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts function performs comparison ‘a’ part >= ‘b’ part. If both inequalities are satisfied, function
returns 1.
Returns
Returns 1 if a >= b, else returns 0.
__device__ unsigned int __vsetgeu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: returns 1 if all 4 pairs compare greater than or equal.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs comparison ‘a’ part >= ‘b’ part. If both inequalities are satisfied, function re-
turns 1.
Returns
Returns 1 if a >= b, else returns 0.
__device__ unsigned int __vsetgts2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: returns 1 if both parts compare greater than.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function performs comparison ‘a’ part > ‘b’ part. If both inequalities are satisfied, function returns
1.
Returns
Returns 1 if a > b, else returns 0.
__device__ unsigned int __vsetgts4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: returns 1 if all 4 pairs compare greater than.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs comparison ‘a’ part > ‘b’ part. If both inequalities are satisfied, function returns
1.
Returns
Returns 1 if a > b, else returns 0.
__device__ unsigned int __vsetgtu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: returns 1 if both parts compare greater than.
14.1. Functions
415
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function performs comparison ‘a’ part > ‘b’ part. If both inequalities are satisfied, function returns
1.
Returns
Returns 1 if a > b, else returns 0.
__device__ unsigned int __vsetgtu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: returns 1 if all 4 pairs compare greater than.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs comparison ‘a’ part > ‘b’ part. If both inequalities are satisfied, function returns
1.
Returns
Returns 1 if a > b, else returns 0.
__device__ unsigned int __vsetles2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: returns 1 if both parts compare less than or equal.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts function performs comparison ‘a’ part <= ‘b’ part. If both inequalities are satisfied, function
returns 1.
Returns
Returns 1 if a <= b, else returns 0.
__device__ unsigned int __vsetles4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: returns 1 if all 4 pairs compare less than or equal.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs comparison ‘a’ part <= ‘b’ part. If both inequalities are satisfied, function re-
turns 1.
Returns
Returns 1 if a <= b, else returns 0.
__device__ unsigned int __vsetleu2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: returns 1 if both parts compare less than or equal.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts function performs comparison ‘a’ part <= ‘b’ part. If both inequalities are satisfied, function
returns 1.
Returns
Returns 1 if a <= b, else returns 0.
__device__ unsigned int __vsetleu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: returns 1 if all 4 pairs compare less than or equal.
Splits 4 bytes of each argument into 4 part, each consisting of 1 byte. For corresponding parts
function performs comparison ‘a’ part <= ‘b’ part. If both inequalities are satisfied, function re-
turns 1.
Returns
Returns 1 if a <= b, else returns 0.
__device__ unsigned int __vsetlts2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: returns 1 if both parts compare less than.
416
Chapter 14. SIMD Intrinsics
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts function performs comparison ‘a’ part <= ‘b’ part. If both inequalities are satisfied, function
returns 1.
Returns
Returns 1 if a < b, else returns 0.
__device__ unsigned int __vsetlts4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: returns 1 if all 4 pairs compare less than.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs comparison ‘a’ part <= ‘b’ part. If both inequalities are satisfied, function re-
turns 1.
Returns
Returns 1 if a < b, else returns 0.
__device__ unsigned int __vsetltu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: returns 1 if both parts compare less than.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts function performs comparison ‘a’ part <= ‘b’ part. If both inequalities are satisfied, function
returns 1.
Returns
Returns 1 if a < b, else returns 0.
__device__ unsigned int __vsetltu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: returns 1 if all 4 pairs compare less than.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs comparison ‘a’ part <= ‘b’ part. If both inequalities are satisfied, function re-
turns 1.
Returns
Returns 1 if a < b, else returns 0.
__device__ unsigned int __vsetne2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed comparison: returns 1 if both parts compare not equal.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function performs comparison ‘a’ part != ‘b’ part. If both conditions are satisfied, function returns
1.
Returns
Returns 1 if a != b, else returns 0.
__device__ unsigned int __vsetne4(unsigned int a, unsigned int b)
Performs per-byte (un)signed comparison: returns 1 if all 4 pairs compare not equal.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs comparison ‘a’ part != ‘b’ part. If both conditions are satisfied, function returns
1.
Returns
Returns 1 if a != b, else returns 0.
__device__ unsigned int __vsub2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed subtraction, with wrap-around: a - b.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding parts
function performs subtraction. Partial results are recombined and returned as unsigned int.
14.1. Functions
417
Returns
Returns computed value.
__device__ unsigned int __vsub4(unsigned int a, unsigned int b)
Performs per-byte subtraction: a - b.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs subtraction. Partial results are recombined and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vsubss2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed subtraction, with signed saturation: a - b.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts function performs subtraction with signed saturation. Partial results are recombined and
returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vsubss4(unsigned int a, unsigned int b)
Performs per-byte subtraction with signed saturation: a - b.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs subtraction with signed saturation. Partial results are recombined and re-
turned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vsubus2(unsigned int a, unsigned int b)
Performs per-halfword subtraction with unsigned saturation: a - b.
Splits 4 bytes of each argument into 2 parts, each consisting of 2 bytes. For corresponding
parts function performs subtraction with unsigned saturation. Partial results are recombined
and returned as unsigned int.
Returns
Returns computed value.
__device__ unsigned int __vsubus4(unsigned int a, unsigned int b)
Performs per-byte subtraction with unsigned saturation: a - b.
Splits 4 bytes of each argument into 4 parts, each consisting of 1 byte. For corresponding parts
function performs subtraction with unsigned saturation. Partial results are recombined and re-
turned as unsigned int.
Returns
Returns computed value.
418
Chapter 14. SIMD Intrinsics
Chapter 15. Structs
15.1. __half
struct __half
__half data type
This structure implements the datatype for storing half-precision floating-point numbers. The
structure implements assignment, arithmetic and comparison operators, and type conversions.
16 bits are being used in total: 1 sign bit, 5 bits for the exponent, and the significand is being
stored in 10 bits. The total precision is 11 bits. There are 15361 representable numbers within the
interval [0.0, 1.0], endpoints included. On average we have log10(2**11) ~ 3.311 decimal digits.
The objective here is to provide IEEE754-compliant implementation of binary16 type and arith-
metic with limitations due to device HW not supporting floating-point exceptions.
Public Functions
__half() = default
Constructor by default.
Emtpy default constructor, result is uninitialized.
__host__ __device__ inline constexpr __half(const __half_raw &hr)
Constructor from __half_raw.
__host__ __device__ explicit __half(const __nv_bfloat16 f)
Construct __half from __nv_bfloat16 input using default round-to-nearest-even round-
ing mode.
Need to include the header file cuda_bf16.h
__host__ __device__ inline __half(const double f)
Construct __half from double input using default round-to-nearest-even rounding mode.
See also:
__double2half(double) for further details.
__host__ __device__ inline __half(const float f)
Construct __half from float input using default round-to-nearest-even rounding mode.
419
See also:
__float2half(float) for further details.
__host__ __device__ inline __half(const int val)
Construct __half from int input using default round-to-nearest-even rounding mode.
__host__ __device__ inline __half(const long long val)
Construct __half from long long input using default round-to-nearest-even rounding
mode.
__host__ __device__ inline __half(const long val)
Construct __half from long input using default round-to-nearest-even rounding mode.
__host__ __device__ inline __half(const short val)
Construct __half from short integer input using default round-to-nearest-even rounding
mode.
__host__ __device__ inline __half(const unsigned int val)
Construct __half from unsigned int input using default round-to-nearest-even rounding
mode.
__host__ __device__ inline __half(const unsigned long long val)
Construct __half from unsigned long long input using default round-to-nearest-even
rounding mode.
__host__ __device__ inline __half(const unsigned long val)
Construct __half from unsigned long input using default round-to-nearest-even round-
ing mode.
__host__ __device__ inline __half(const unsigned short val)
Construct __half from unsigned short integer input using default round-to-nearest-even
rounding mode.
__host__ __device__ operator
__half_raw() const
Type cast to __half_raw operator.
__host__ __device__ operator
__half_raw() volatile const
Type cast to __half_raw operator with volatile input.
__host__ __device__ inline constexpr operator bool() const
Conversion operator to bool data type.
+0 and -0 inputs convert to false. Non-zero inputs convert to true.
__host__ __device__ operator char() const
Conversion operator to an implementation defined char data type.
Using round-toward-zero rounding mode.
Detects signedness of the char type and proceeds accordingly, see further details in
__half2char_rz(__half) and __half2uchar_rz(__half).
__host__ __device__ operator float() const
Type cast to float operator.
420
Chapter 15. Structs
__host__ __device__ operator int() const
Conversion operator to int data type.
Using round-toward-zero rounding mode.
See also:
__half2int_rz(__half) for further details.
__host__ __device__ operator long() const
Conversion operator to long data type.
Using round-toward-zero rounding mode.
Detects size of the long type and proceeds accordingly, see further details in
__half2int_rz(__half) and __half2ll_rz(__half).
__host__ __device__ operator long long() const
Conversion operator to long long data type.
Using round-toward-zero rounding mode.
See also:
__half2ll_rz(__half) for further details.
__host__ __device__ operator short() const
Conversion operator to short data type.
Using round-toward-zero rounding mode.
See also:
__half2short_rz(__half) for further details.
__host__ __device__ operator signed char() const
Conversion operator to signed char data type.
Using round-toward-zero rounding mode.
See also:
__half2char_rz(__half) for further details.
__host__ __device__ operator unsigned char() const
Conversion operator to unsigned char data type.
Using round-toward-zero rounding mode.
See also:
__half2uchar_rz(__half) for further details.
15.1.
__half
421
__host__ __device__ operator unsigned int() const
Conversion operator to unsigned int data type.
Using round-toward-zero rounding mode.
See also:
__half2uint_rz(__half) for further details.
__host__ __device__ operator unsigned long() const
Conversion operator to unsigned long data type.
Using round-toward-zero rounding mode.
Detects size of the unsigned long type and proceeds accordingly, see further details in
__half2uint_rz(__half) and __half2ull_rz(__half).
__host__ __device__ operator unsigned long long() const
Conversion operator to unsigned long long data type.
Using round-toward-zero rounding mode.
See also:
__half2ull_rz(__half) for further details.
__host__ __device__ operator unsigned short() const
Conversion operator to unsigned short data type.
Using round-toward-zero rounding mode.
See also:
__half2ushort_rz(__half) for further details.
__host__ __device__ __half &operator=(const __half_raw &hr)
Assignment operator from __half_raw.
__host__ __device__ volatile __half &operator=(const __half_raw &hr) volatile
Assignment operator from __half_raw to volatile __half.
__host__ __device__ __half &operator=(const double f)
Type cast to __half assignment operator from double input using default round-to-
nearest-even rounding mode.
See also:
__double2half(double) for further details.
__host__ __device__ __half &operator=(const float f)
Type cast to __half assignment operator from float input using default round-to-nearest-
even rounding mode.
See also:
__float2half(float) for further details.
422
Chapter 15. Structs
__host__ __device__ __half &operator=(const int val)
Type cast from int assignment operator, using default round-to-nearest-even rounding
mode.
__host__ __device__ __half &operator=(const long long val)
Type cast from long long assignment operator, using default round-to-nearest-even
rounding mode.
__host__ __device__ __half &operator=(const short val)
Type cast from short assignment operator, using default round-to-nearest-even rounding
mode.
__host__ __device__ __half &operator=(const unsigned int val)
Type cast from unsigned int assignment operator, using default round-to-nearest-even
rounding mode.
__host__ __device__ __half &operator=(const unsigned long long val)
Type cast from unsigned long long assignment operator, using default round-to-nearest-
even rounding mode.
__host__ __device__ __half &operator=(const unsigned short val)
Type cast from unsigned short assignment operator, using default round-to-nearest-even
rounding mode.
__host__ __device__ volatile __half &operator=(volatile const __half_raw &hr) volatile
Assignment operator from volatile __half_raw to volatile __half.
15.2. __half2
struct __half2
__half2 data type
This structure implements the datatype for storing two half-precision floating-point numbers.
The structure implements assignment, arithmetic and comparison operators, and type conver-
sions.
▶ NOTE: __half2 is visible to non-nvcc host compilers
Public Functions
__half2() = default
Constructor by default.
Emtpy default constructor, result is uninitialized.
__host__ __device__ inline constexpr __half2(const __half &a, const __half &b)
Constructor from two __half variables.
__host__ __device__ inline __half2(const __half2 &&src)
Move constructor, available for C++11 and later dialects.
15.2.
__half2
423
__host__ __device__ inline __half2(const __half2 &src)
Copy constructor.
__host__ __device__ inline __half2(const __half2_raw &h2r)
Constructor from __half2_raw.
__host__ __device__ operator
__half2_raw() const
Conversion operator to __half2_raw.
__host__ __device__ __half2 &operator=(const __half2 &&src)
Move assignment operator, available for C++11 and later dialects.
__host__ __device__ __half2 &operator=(const __half2 &src)
Copy assignment operator.
__host__ __device__ __half2 &operator=(const __half2_raw &h2r)
Assignment operator from __half2_raw.
Public Members
__half x
Storage field holding lower __half part.
__half y
Storage field holding upper __half part.
15.3. __half2_raw
struct __half2_raw
__half2_raw data type
Type allows static initialization of half2 until it becomes a built-in type.
▶ Note: this initialization is as a bit-field representation of half2, and not a conversion from
short2 to half2. Such representation will be deprecated in a future version of CUDA.
▶ Note: this is visible to non-nvcc compilers, including C-only compilations
Public Members
unsigned short x
Storage field contains bits of the lower half part.
unsigned short y
Storage field contains bits of the upper half part.
424
Chapter 15. Structs
15.4. __half_raw
struct __half_raw
__half_raw data type
Type allows static initialization of half until it becomes a built-in type.
▶ Note: this initialization is as a bit-field representation of half, and not a conversion from
short to half. Such representation will be deprecated in a future version of CUDA.
▶ Note: this is visible to non-nvcc compilers, including C-only compilations
Public Members
unsigned short x
Storage field contains bits representation of the half floating-point number.
15.5. __nv_bfloat16
struct __nv_bfloat16
nv_bfloat16 datatype
This structure implements the datatype for storing nv_bfloat16 floating-point numbers. The
structure implements assignment operators and type conversions. 16 bits are being used in
total: 1 sign bit, 8 bits for the exponent, and the significand is being stored in 7 bits. The total
precision is 8 bits.
Public Functions
__nv_bfloat16() = default
Constructor by default.
Emtpy default constructor, result is uninitialized.
__host__ __device__ inline explicit __nv_bfloat16(const __half f)
Construct __nv_bfloat16 from __half input using default round-to-nearest-even round-
ing mode.
__host__ __device__ inline constexpr __nv_bfloat16(const __nv_bfloat16_raw &hr)
Constructor from __nv_bfloat16_raw.
__host__ __device__ inline __nv_bfloat16(const double f)
Construct __nv_bfloat16 from double input using default round-to-nearest-even round-
ing mode.
15.4.
__half_raw
425
__host__ __device__ inline __nv_bfloat16(const float f)
Construct __nv_bfloat16 from float input using default round-to-nearest-even rounding
mode.
__host__ __device__ inline __nv_bfloat16(const long val)
Construct __nv_bfloat16 from long input using default round-to-nearest-even rounding
mode.
__host__ __device__ inline __nv_bfloat16(const unsigned long val)
Construct __nv_bfloat16 from unsigned long input using default round-to-nearest-even
rounding mode.
__host__ __device__ inline __nv_bfloat16(int val)
Construct __nv_bfloat16 from int input using default round-to-nearest-even rounding
mode.
__host__ __device__ inline __nv_bfloat16(long long val)
Construct __nv_bfloat16 from long long input using default round-to-nearest-even
rounding mode.
__host__ __device__ inline __nv_bfloat16(short val)
Construct __nv_bfloat16 from short integer input using default round-to-nearest-even
rounding mode.
__host__ __device__ inline __nv_bfloat16(unsigned int val)
Construct __nv_bfloat16 from unsigned int input using default round-to-nearest-even
rounding mode.
__host__ __device__ inline __nv_bfloat16(unsigned long long val)
Construct __nv_bfloat16 from unsigned long long input using default round-to-
nearest-even rounding mode.
__host__ __device__ inline __nv_bfloat16(unsigned short val)
Construct __nv_bfloat16 from unsigned short integer input using default round-to-
nearest-even rounding mode.
__host__ __device__ operator
__nv_bfloat16_raw() const
Type cast to __nv_bfloat16_raw operator.
__host__ __device__ operator
__nv_bfloat16_raw() volatile const
Type cast to __nv_bfloat16_raw operator with volatile input.
__host__ __device__ inline constexpr operator bool() const
Conversion operator to bool data type.
+0 and -0 inputs convert to false. Non-zero inputs convert to true.
__host__ __device__ operator char() const
Conversion operator to an implementation defined char data type.
Using round-toward-zero rounding mode.
Detects signedness of the char type and proceeds accordingly, see further details in signed
and unsigned char operators.
__host__ __device__ operator float() const
Type cast to float operator.
426
Chapter 15. Structs
__host__ __device__ operator int() const
Conversion operator to int data type.
Using round-toward-zero rounding mode.
See __bfloat162int_rz(__nv_bfloat16) for further details
__host__ __device__ operator long() const
Conversion operator to long data type.
Using round-toward-zero rounding mode.
__host__ __device__ operator long long() const
Conversion operator to long long data type.
Using round-toward-zero rounding mode.
See __bfloat162ll_rz(__nv_bfloat16) for further details
__host__ __device__ operator short() const
Conversion operator to short data type.
Using round-toward-zero rounding mode.
See __bfloat162short_rz(__nv_bfloat16) for further details
__host__ __device__ operator signed char() const
Conversion operator to signed char data type.
Using round-toward-zero rounding mode.
See __bfloat162char_rz(__nv_bfloat16) for further details
__host__ __device__ operator unsigned char() const
Conversion operator to unsigned char data type.
Using round-toward-zero rounding mode.
See __bfloat162uchar_rz(__nv_bfloat16) for further details
__host__ __device__ operator unsigned int() const
Conversion operator to unsigned int data type.
Using round-toward-zero rounding mode.
See __bfloat162uint_rz(__nv_bfloat16) for further details
__host__ __device__ operator unsigned long() const
Conversion operator to unsigned long data type.
Using round-toward-zero rounding mode.
__host__ __device__ operator unsigned long long() const
Conversion operator to unsigned long long data type.
Using round-toward-zero rounding mode.
See __bfloat162ull_rz(__nv_bfloat16) for further details
__host__ __device__ operator unsigned short() const
Conversion operator to unsigned short data type.
Using round-toward-zero rounding mode.
See __bfloat162ushort_rz(__nv_bfloat16) for further details
15.5.
__nv_bfloat16
427
__host__ __device__ __nv_bfloat16 &operator=(const __nv_bfloat16_raw &hr)
Assignment operator from __nv_bfloat16_raw.
__host__ __device__ volatile __nv_bfloat16 &operator=(const __nv_bfloat16_raw &hr) volatile
Assignment operator from __nv_bfloat16_raw to volatile __nv_bfloat16.
__host__ __device__ __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 &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 &operator=(volatile const __nv_bfloat16_raw &hr)
volatile
Assignment operator from volatile __nv_bfloat16_raw to volatile __nv_bfloat16.
__host__ __device__ __nv_bfloat16 &operator=(int val)
Type cast from int assignment operator, using default round-to-nearest-even rounding
mode.
__host__ __device__ __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 &operator=(short val)
Type cast from short assignment operator, using default round-to-nearest-even rounding
mode.
__host__ __device__ __nv_bfloat16 &operator=(unsigned int val)
Type cast from unsigned int assignment operator, using default round-to-nearest-even
rounding mode.
__host__ __device__ __nv_bfloat16 &operator=(unsigned long long val)
Type cast from unsigned long long assignment operator, using default round-to-nearest-
even rounding mode.
__host__ __device__ __nv_bfloat16 &operator=(unsigned short val)
Type cast from unsigned short assignment operator, using default round-to-nearest-even
rounding mode.
15.6. __nv_bfloat162
struct __nv_bfloat162
nv_bfloat162 datatype
This structure implements the datatype for storing two nv_bfloat16 floating-point numbers. The
structure implements assignment, arithmetic and comparison operators, and type conversions.
▶ NOTE: __nv_bfloat162 is visible to non-nvcc host compilers
428
Chapter 15. Structs
Public Functions
__nv_bfloat162() = default
Constructor by default.
Emtpy default constructor, result is uninitialized.
__host__ __device__ __nv_bfloat162(__nv_bfloat162 &&src)
Move constructor, available for C++11 and later dialects.
__host__ __device__ inline constexpr __nv_bfloat162(const __nv_bfloat16 &a, const
__nv_bfloat16 &b)
Constructor from two __nv_bfloat16 variables.
__host__ __device__ __nv_bfloat162(const __nv_bfloat162 &src)
Copy constructor.
__host__ __device__ __nv_bfloat162(const __nv_bfloat162_raw &h2r)
Constructor from __nv_bfloat162_raw.
__host__ __device__ operator
__nv_bfloat162_raw() const
Conversion operator to __nv_bfloat162_raw.
__host__ __device__ __nv_bfloat162 &operator=(__nv_bfloat162 &&src)
Move assignment operator, available for C++11 and later dialects.
__host__ __device__ __nv_bfloat162 &operator=(const __nv_bfloat162 &src)
Copy assignment operator.
__host__ __device__ __nv_bfloat162 &operator=(const __nv_bfloat162_raw &h2r)
Assignment operator from __nv_bfloat162_raw.
Public Members
__nv_bfloat16 x
Storage field holding lower __nv_bfloat16 part.
__nv_bfloat16 y
Storage field holding upper __nv_bfloat16 part.
15.7. __nv_bfloat162_raw
struct __nv_bfloat162_raw
__nv_bfloat162_raw data type
Type allows static initialization of nv_bfloat162 until it becomes a built-in type.
▶ Note: this initialization is as a bit-field representation of nv_bfloat162, and not a conver-
sion from short2 to nv_bfloat162. Such representation will be deprecated in a future
version of CUDA.
15.7.
__nv_bfloat162_raw
429
▶ Note: this is visible to non-nvcc compilers, including C-only compilations
Public Members
unsigned short x
Storage field contains bits of the lower nv_bfloat16 part.
unsigned short y
Storage field contains bits of the upper nv_bfloat16 part.
15.8. __nv_bfloat16_raw
struct __nv_bfloat16_raw
__nv_bfloat16_raw data type
Type allows static initialization of nv_bfloat16 until it becomes a built-in type.
▶ Note: this initialization is as a bit-field representation of nv_bfloat16, and not a conversion
from short to nv_bfloat16. Such representation will be deprecated in a future version of
CUDA.
▶ Note: this is visible to non-nvcc compilers, including C-only compilations
Public Members
unsigned short x
Storage field contains bits representation of the nv_bfloat16 floating-point number.
15.9. __nv_fp4_e2m1
struct __nv_fp4_e2m1
__nv_fp4_e2m1 datatype
This structure implements the datatype for handling fp4 floating-point numbers of e2m1 kind:
with 1 sign, 2 exponent, 1 implicit and 1 explicit mantissa bits. This encoding does not support
Inf/NaN.
The structure implements converting constructors and operators.
430
Chapter 15. Structs
Public Functions
__host__ __device__ inline __nv_fp4_e2m1()
Constructor by default.
__host__ __device__ inline explicit __nv_fp4_e2m1(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for out-of-range
values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp4_e2m1(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for out-
of-range values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp4_e2m1(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for out-of-range
values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp4_e2m1(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior for out-of-range
values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp4_e2m1(const int val)
Constructor from int data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ inline explicit __nv_fp4_e2m1(const long int val)
Constructor from long int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp4_e2m1(const long long int val)
Constructor from long long int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp4_e2m1(const short int val)
Constructor from short int data type.
__host__ __device__ inline explicit __nv_fp4_e2m1(const unsigned int val)
Constructor from unsigned int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp4_e2m1(const unsigned long int val)
Constructor from unsigned long int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ inline explicit __nv_fp4_e2m1(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on __NV_SATFINITE behavior
for out-of-range values.
__host__ __device__ inline explicit __nv_fp4_e2m1(const unsigned short int val)
Constructor from unsigned short int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
15.9.
__nv_fp4_e2m1
431
Public Members
__nv_fp4_storage_t __x
Storage variable contains the fp4 floating-point data.
15.10. __nv_fp4x2_e2m1
struct __nv_fp4x2_e2m1
__nv_fp4x2_e2m1 datatype
This structure implements the datatype for handling two fp4 floating-point numbers of e2m1
kind each.
The structure implements converting constructors and operators.
Public Functions
__host__ __device__ inline __nv_fp4x2_e2m1()
Constructor by default.
__host__ __device__ inline explicit __nv_fp4x2_e2m1(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp4x2_e2m1(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp4x2_e2m1(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp4x2_e2m1(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
Public Members
__nv_fp4x2_storage_t __x
Storage variable contains the vector of two fp4 floating-point data values.
432
Chapter 15. Structs
15.11. __nv_fp4x4_e2m1
struct __nv_fp4x4_e2m1
__nv_fp4x4_e2m1 datatype
This structure implements the datatype for handling four fp4 floating-point numbers of e2m1
kind each.
The structure implements converting constructors and operators.
Public Functions
__host__ __device__ inline __nv_fp4x4_e2m1()
Constructor by default.
__host__ __device__ inline explicit __nv_fp4x4_e2m1(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior
for out-of-range values.
__host__ __device__ inline explicit __nv_fp4x4_e2m1(const __nv_bfloat162 flo, const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
__host__ __device__ inline explicit __nv_fp4x4_e2m1(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp4x4_e2m1(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
Public Members
__nv_fp4x4_storage_t __x
Storage variable contains the vector of four fp4 floating-point data values.
15.12. __nv_fp6_e2m3
struct __nv_fp6_e2m3
__nv_fp6_e2m3 datatype
This structure implements the datatype for storing fp6 floating-point numbers of e2m3 kind:
with 1 sign, 2 exponent, 1 implicit and 3 explicit mantissa bits. This encoding does not support
Inf/NaN.
The structure implements converting constructors and operators.
15.11. __nv_fp4x4_e2m1
433
Public Functions
__host__ __device__ inline __nv_fp6_e2m3()
Constructor by default.
__host__ __device__ inline explicit __nv_fp6_e2m3(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for out-of-range
values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp6_e2m3(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for out-
of-range values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp6_e2m3(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for out-of-range
values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp6_e2m3(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior for out-of-range
values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp6_e2m3(const int val)
Constructor from int data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ inline explicit __nv_fp6_e2m3(const long int val)
Constructor from long int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp6_e2m3(const long long int val)
Constructor from long long int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp6_e2m3(const short int val)
Constructor from short int data type.
__host__ __device__ inline explicit __nv_fp6_e2m3(const unsigned int val)
Constructor from unsigned int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp6_e2m3(const unsigned long int val)
Constructor from unsigned long int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ inline explicit __nv_fp6_e2m3(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on __NV_SATFINITE behavior
for out-of-range values.
__host__ __device__ inline explicit __nv_fp6_e2m3(const unsigned short int val)
Constructor from unsigned short int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
434
Chapter 15. Structs
Public Members
__nv_fp6_storage_t __x
Storage variable contains the fp6 floating-point data.
15.13. __nv_fp6_e3m2
struct __nv_fp6_e3m2
__nv_fp6_e3m2 datatype
This structure implements the datatype for handling fp6 floating-point numbers of e3m2 kind:
with 1 sign, 3 exponent, 1 implicit and 2 explicit mantissa bits. This encoding does not support
Inf/NaN.
The structure implements converting constructors and operators.
Public Functions
__host__ __device__ inline __nv_fp6_e3m2()
Constructor by default.
__host__ __device__ inline explicit __nv_fp6_e3m2(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for out-of-range
values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp6_e3m2(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for out-
of-range values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp6_e3m2(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for out-of-range
values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp6_e3m2(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior for out-of-range
values and cudaRoundNearest rounding mode.
__host__ __device__ inline explicit __nv_fp6_e3m2(const int val)
Constructor from int data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ inline explicit __nv_fp6_e3m2(const long int val)
Constructor from long int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp6_e3m2(const long long int val)
Constructor from long long int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp6_e3m2(const short int val)
Constructor from short int data type.
15.13. __nv_fp6_e3m2
435
__host__ __device__ inline explicit __nv_fp6_e3m2(const unsigned int val)
Constructor from unsigned int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp6_e3m2(const unsigned long int val)
Constructor from unsigned long int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ inline explicit __nv_fp6_e3m2(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on __NV_SATFINITE behavior
for out-of-range values.
__host__ __device__ inline explicit __nv_fp6_e3m2(const unsigned short int val)
Constructor from unsigned short int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
Public Members
__nv_fp6_storage_t __x
Storage variable contains the fp6 floating-point data.
15.14. __nv_fp6x2_e2m3
struct __nv_fp6x2_e2m3
__nv_fp6x2_e2m3 datatype
This structure implements the datatype for handling two fp6 floating-point numbers of e2m3
kind each.
The structure implements converting constructors and operators.
Public Functions
__host__ __device__ inline __nv_fp6x2_e2m3()
Constructor by default.
__host__ __device__ inline explicit __nv_fp6x2_e2m3(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp6x2_e2m3(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp6x2_e2m3(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp6x2_e2m3(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
436
Chapter 15. Structs
Public Members
__nv_fp6x2_storage_t __x
Storage variable contains the vector of two fp6 floating-point data values.
15.15. __nv_fp6x2_e3m2
struct __nv_fp6x2_e3m2
__nv_fp6x2_e3m2 datatype
This structure implements the datatype for handling two fp6 floating-point numbers of e3m2
kind each.
The structure implements converting constructors and operators.
Public Functions
__host__ __device__ inline __nv_fp6x2_e3m2()
Constructor by default.
__host__ __device__ inline explicit __nv_fp6x2_e3m2(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp6x2_e3m2(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp6x2_e3m2(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp6x2_e3m2(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
Public Members
__nv_fp6x2_storage_t __x
Storage variable contains the vector of two fp6 floating-point data values.
15.15. __nv_fp6x2_e3m2
437
15.16. __nv_fp6x4_e2m3
struct __nv_fp6x4_e2m3
__nv_fp6x4_e2m3 datatype
This structure implements the datatype for handling four fp6 floating-point numbers of e2m3
kind each.
The structure implements converting constructors and operators.
Public Functions
__host__ __device__ inline __nv_fp6x4_e2m3()
Constructor by default.
__host__ __device__ inline explicit __nv_fp6x4_e2m3(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior
for out-of-range values.
__host__ __device__ inline explicit __nv_fp6x4_e2m3(const __nv_bfloat162 flo, const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
__host__ __device__ inline explicit __nv_fp6x4_e2m3(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp6x4_e2m3(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
Public Members
__nv_fp6x4_storage_t __x
Storage variable contains the vector of four fp6 floating-point data values.
15.17. __nv_fp6x4_e3m2
struct __nv_fp6x4_e3m2
__nv_fp6x4_e3m2 datatype
This structure implements the datatype for handling four fp6 floating-point numbers of e3m2
kind each.
The structure implements converting constructors and operators.
438
Chapter 15. Structs
Public Functions
__host__ __device__ inline __nv_fp6x4_e3m2()
Constructor by default.
__host__ __device__ inline explicit __nv_fp6x4_e3m2(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior
for out-of-range values.
__host__ __device__ inline explicit __nv_fp6x4_e3m2(const __nv_bfloat162 flo, const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
__host__ __device__ inline explicit __nv_fp6x4_e3m2(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp6x4_e3m2(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
Public Members
__nv_fp6x4_storage_t __x
Storage variable contains the vector of four fp6 floating-point data values.
15.18. __nv_fp8_e4m3
struct __nv_fp8_e4m3
__nv_fp8_e4m3 datatype
This structure implements the datatype for storing fp8 floating-point numbers of e4m3 kind:
with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. The encoding doesn’t support
Infinity. NaNs are limited to 0x7F and 0xFF values.
The structure implements converting constructors and operators.
Public Functions
__nv_fp8_e4m3() = default
Constructor by default.
__host__ __device__ inline explicit __nv_fp8_e4m3(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for out-of-range
values.
15.18. __nv_fp8_e4m3
439
__host__ __device__ inline explicit __nv_fp8_e4m3(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp8_e4m3(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8_e4m3(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8_e4m3(const int val)
Constructor from int data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ inline explicit __nv_fp8_e4m3(const long int val)
Constructor from long int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8_e4m3(const long long int val)
Constructor from long long int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp8_e4m3(const short int val)
Constructor from short int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp8_e4m3(const unsigned int val)
Constructor from unsigned int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp8_e4m3(const unsigned long int val)
Constructor from unsigned long int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ inline explicit __nv_fp8_e4m3(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on __NV_SATFINITE behavior
for out-of-range values.
__host__ __device__ inline explicit __nv_fp8_e4m3(const unsigned short int val)
Constructor from unsigned short int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ inline explicit operator
__half() const
Conversion operator to __half data type.
__host__ __device__ inline explicit operator
__nv_bfloat16() const
Conversion operator to __nv_bfloat16 data type.
__host__ __device__ inline explicit operator bool() const
Conversion operator to bool data type.
+0 and -0 inputs convert to false. Non-zero inputs convert to true.
440
Chapter 15. Structs
__host__ __device__ inline explicit operator char() const
Conversion operator to an implementation defined char data type.
Detects signedness of the char type and proceeds accordingly, see further details in signed
and unsigned char operators.
Clamps inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator double() const
Conversion operator to double data type.
__host__ __device__ inline explicit operator float() const
Conversion operator to float data type.
__host__ __device__ inline explicit operator int() const
Conversion operator to int data type.
NaN inputs convert to zero.
__host__ __device__ inline explicit operator long int() const
Conversion operator to long int data type.
Clamps too large inputs to the output range. NaN inputs convert to zero if output type is
32-bit. NaN inputs convert to 0x8000000000000000ULL if output type is 64-bit.
__host__ __device__ inline explicit operator long long int() const
Conversion operator to long long int data type.
NaN inputs convert to 0x8000000000000000LL.
__host__ __device__ inline explicit operator short int() const
Conversion operator to short int data type.
NaN inputs convert to zero.
__host__ __device__ inline explicit operator signed char() const
Conversion operator to signed char data type.
Clamps too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator unsigned char() const
Conversion operator to unsigned char data type.
Clamps negative and too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator unsigned int() const
Conversion operator to unsigned int data type.
Clamps negative inputs to zero. NaN inputs convert to zero.
__host__ __device__ inline explicit operator unsigned long int() const
Conversion operator to unsigned long int data type.
Clamps negative and too large inputs to the output range. NaN inputs convert to zero if
output type is 32-bit. NaN inputs convert to 0x8000000000000000ULL if output type is
64-bit.
__host__ __device__ inline explicit operator unsigned long long int() const
Conversion operator to unsigned long long int data type.
Clamps negative inputs to zero. NaN inputs convert to 0x8000000000000000ULL.
15.18. __nv_fp8_e4m3
441
__host__ __device__ inline explicit operator unsigned short int() const
Conversion operator to unsigned short int data type.
Clamps negative inputs to zero. NaN inputs convert to zero.
Public Members
__nv_fp8_storage_t __x
Storage variable contains the fp8 floating-point data.
15.19. __nv_fp8_e5m2
struct __nv_fp8_e5m2
__nv_fp8_e5m2 datatype
This structure implements the datatype for handling fp8 floating-point numbers of e5m2 kind:
with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits.
The structure implements converting constructors and operators.
Public Functions
__nv_fp8_e5m2() = default
Constructor by default.
__host__ __device__ inline explicit __nv_fp8_e5m2(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8_e5m2(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp8_e5m2(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8_e5m2(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8_e5m2(const int val)
Constructor from int data type, relies on __NV_SATFINITE behavior for out-of-range val-
ues.
__host__ __device__ inline explicit __nv_fp8_e5m2(const long int val)
Constructor from long int data type, relies on __NV_SATFINITE behavior for out-of-range
values.
442
Chapter 15. Structs
__host__ __device__ inline explicit __nv_fp8_e5m2(const long long int val)
Constructor from long long int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp8_e5m2(const short int val)
Constructor from short int data type.
__host__ __device__ inline explicit __nv_fp8_e5m2(const unsigned int val)
Constructor from unsigned int data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit __nv_fp8_e5m2(const unsigned long int val)
Constructor from unsigned long int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ inline explicit __nv_fp8_e5m2(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on __NV_SATFINITE behavior
for out-of-range values.
__host__ __device__ inline explicit __nv_fp8_e5m2(const unsigned short int val)
Constructor from unsigned short int data type, relies on __NV_SATFINITE behavior for
out-of-range values.
__host__ __device__ inline explicit operator
__half() const
Conversion operator to __half data type.
__host__ __device__ inline explicit operator
__nv_bfloat16() const
Conversion operator to __nv_bfloat16 data type.
__host__ __device__ inline explicit operator bool() const
Conversion operator to bool data type.
+0 and -0 inputs convert to false. Non-zero inputs convert to true.
__host__ __device__ inline explicit operator char() const
Conversion operator to an implementation defined char data type.
Detects signedness of the char type and proceeds accordingly, see further details in signed
and unsigned char operators.
Clamps inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator double() const
Conversion operator to double data type.
__host__ __device__ inline explicit operator float() const
Conversion operator to float data type.
__host__ __device__ inline explicit operator int() const
Conversion operator to int data type.
Clamps too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator long int() const
Conversion operator to long int data type.
Clamps too large inputs to the output range. NaN inputs convert to zero if output type is
32-bit. NaN inputs convert to 0x8000000000000000ULL if output type is 64-bit.
15.19. __nv_fp8_e5m2
443
__host__ __device__ inline explicit operator long long int() const
Conversion operator to long long int data type.
Clamps too large inputs to the output range.
NaN inputs convert to
0x8000000000000000LL.
__host__ __device__ inline explicit operator short int() const
Conversion operator to short int data type.
Clamps too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator signed char() const
Conversion operator to signed char data type.
Clamps too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator unsigned char() const
Conversion operator to unsigned char data type.
Clamps negative and too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator unsigned int() const
Conversion operator to unsigned int data type.
Clamps negative and too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator unsigned long int() const
Conversion operator to unsigned long int data type.
Clamps negative and too large inputs to the output range. NaN inputs convert to zero if
output type is 32-bit. NaN inputs convert to 0x8000000000000000ULL if output type is
64-bit.
__host__ __device__ inline explicit operator unsigned long long int() const
Conversion operator to unsigned long long int data type.
Clamps negative and too large inputs to the output range.
NaN inputs convert to
0x8000000000000000ULL.
__host__ __device__ inline explicit operator unsigned short int() const
Conversion operator to unsigned short int data type.
Clamps negative and too large inputs to the output range. NaN inputs convert to zero.
Public Members
__nv_fp8_storage_t __x
Storage variable contains the fp8 floating-point data.
444
Chapter 15. Structs
15.20. __nv_fp8_e8m0
struct __nv_fp8_e8m0
__nv_fp8_e8m0 datatype
This structure implements the datatype for handling 8-bit scale factors of e8m0 kind: interpreted
as powers of two with biased exponent. Bias equals to 127, so numbers 0 through 254 represent
2^-127 through 2^127. Number 0xFF = 255 is reserved for NaN.
The structure implements converting constructors and operators.
Public Functions
__nv_fp8_e8m0() = default
Constructor by default.
__host__ __device__ inline explicit __nv_fp8_e8m0(const __half f)
Constructor from __half data type, relies on __NV_SATFINITE behavior for large input
values and cudaRoundPosInf for rounding.
See also:
__nv_cvt_float_to_e8m0 for further details
__host__ __device__ inline explicit __nv_fp8_e8m0(const __nv_bfloat16 f)
Constructor from __nv_bfloat16 data type, relies on __NV_SATFINITE behavior for large
input values and cudaRoundPosInf for rounding.
See also:
__nv_cvt_bfloat16raw_to_e8m0 for further details
__host__ __device__ inline explicit __nv_fp8_e8m0(const double f)
Constructor from double data type, relies on __NV_SATFINITE behavior for large input
values and cudaRoundPosInf for rounding.
See also:
__nv_cvt_double_to_e8m0 for further details
__host__ __device__ inline explicit __nv_fp8_e8m0(const float f)
Constructor from float data type, relies on __NV_SATFINITE behavior behavior for large
input values and cudaRoundPosInf for rounding.
See also:
__nv_cvt_float_to_e8m0 for further details
__host__ __device__ inline explicit __nv_fp8_e8m0(const int val)
Constructor from int data type, relies on cudaRoundPosInf rounding.
15.20. __nv_fp8_e8m0
445
__host__ __device__ inline explicit __nv_fp8_e8m0(const long int val)
Constructor from long int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ inline explicit __nv_fp8_e8m0(const long long int val)
Constructor from long long int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ inline explicit __nv_fp8_e8m0(const short int val)
Constructor from short int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ inline explicit __nv_fp8_e8m0(const unsigned int val)
Constructor from unsigned int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ inline explicit __nv_fp8_e8m0(const unsigned long int val)
Constructor from unsigned long int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ inline explicit __nv_fp8_e8m0(const unsigned long long int val)
Constructor from unsigned long long int data type, relies on cudaRoundPosInf round-
ing.
__host__ __device__ inline explicit __nv_fp8_e8m0(const unsigned short int val)
Constructor from unsigned short int data type, relies on cudaRoundPosInf rounding.
__host__ __device__ inline explicit operator
__half() const
Conversion operator to __half data type.
__host__ __device__ inline explicit operator
__nv_bfloat16() const
Conversion operator to __nv_bfloat16 data type.
__host__ __device__ inline explicit operator bool() const
Conversion operator to bool data type.
All values in input range are non-zero, so result is always true.
__host__ __device__ inline explicit operator char() const
Conversion operator to an implementation defined char data type.
Detects signedness of the char type and proceeds accordingly, see further details in signed
and unsigned char operators.
Clamps inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator double() const
Conversion operator to double data type.
__host__ __device__ inline explicit operator float() const
Conversion operator to float data type.
__host__ __device__ inline explicit operator int() const
Conversion operator to int data type.
Clamps too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator long int() const
Conversion operator to long int data type.
Clamps too large inputs to the output range. NaN inputs convert to zero if output type is
32-bit. NaN inputs convert to 0x8000000000000000ULL if output type is 64-bit.
446
Chapter 15. Structs
__host__ __device__ inline explicit operator long long int() const
Conversion operator to long long int data type.
Clamps too large inputs to the output range.
NaN inputs convert to
0x8000000000000000LL.
__host__ __device__ inline explicit operator short int() const
Conversion operator to short int data type.
Clamps too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator signed char() const
Conversion operator to signed char data type.
Clamps too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator unsigned char() const
Conversion operator to unsigned char data type.
Clamps too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator unsigned int() const
Conversion operator to unsigned int data type.
Clamps too large inputs to the output range. NaN inputs convert to zero.
__host__ __device__ inline explicit operator unsigned long int() const
Conversion operator to unsigned long int data type.
Clamps too large inputs to the output range. NaN inputs convert to zero if output type is
32-bit. NaN inputs convert to 0x8000000000000000ULL if output type is 64-bit.
__host__ __device__ inline explicit operator unsigned long long int() const
Conversion operator to unsigned long long int data type.
Clamps too large inputs to the output range.
NaN inputs convert to
0x8000000000000000ULL.
__host__ __device__ inline explicit operator unsigned short int() const
Conversion operator to unsigned short int data type.
Clamps too large inputs to the output range. NaN inputs convert to zero.
Public Members
__nv_fp8_storage_t __x
Storage variable contains the 8-bit scale data.
15.20. __nv_fp8_e8m0
447
15.21. __nv_fp8x2_e4m3
struct __nv_fp8x2_e4m3
__nv_fp8x2_e4m3 datatype
This structure implements the datatype for storage and operations on the vector of two fp8
values of e4m3 kind each: with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. The
encoding doesn’t support Infinity. NaNs are limited to 0x7F and 0xFF values.
Public Functions
__nv_fp8x2_e4m3() = default
Constructor by default.
__host__ __device__ inline explicit __nv_fp8x2_e4m3(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8x2_e4m3(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp8x2_e4m3(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8x2_e4m3(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit operator
__half2() const
Conversion operator to __half2 data type.
__host__ __device__ inline explicit operator float2() const
Conversion operator to float2 data type.
Public Members
__nv_fp8x2_storage_t __x
Storage variable contains the vector of two fp8 floating-point data values.
448
Chapter 15. Structs
15.22. __nv_fp8x2_e5m2
struct __nv_fp8x2_e5m2
__nv_fp8x2_e5m2 datatype
This structure implements the datatype for handling two fp8 floating-point numbers of e5m2
kind each: with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits.
The structure implements converting constructors and operators.
Public Functions
__nv_fp8x2_e5m2() = default
Constructor by default.
__host__ __device__ inline explicit __nv_fp8x2_e5m2(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8x2_e5m2(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp8x2_e5m2(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8x2_e5m2(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit operator
__half2() const
Conversion operator to __half2 data type.
__host__ __device__ inline explicit operator float2() const
Conversion operator to float2 data type.
Public Members
__nv_fp8x2_storage_t __x
Storage variable contains the vector of two fp8 floating-point data values.
15.22. __nv_fp8x2_e5m2
449
15.23. __nv_fp8x2_e8m0
struct __nv_fp8x2_e8m0
__nv_fp8x2_e8m0 datatype
This structure implements the datatype for storage and operations on the vector of two scale
factors of e8m0 kind each.
Public Functions
__nv_fp8x2_e8m0() = default
Constructor by default.
__host__ __device__ inline explicit __nv_fp8x2_e8m0(const __half2 f)
Constructor from __half2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8x2_e8m0(const __nv_bfloat162 f)
Constructor from __nv_bfloat162 data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp8x2_e8m0(const double2 f)
Constructor from double2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit __nv_fp8x2_e8m0(const float2 f)
Constructor from float2 data type, relies on __NV_SATFINITE behavior for out-of-range
values.
__host__ __device__ inline explicit operator
__half2() const
Conversion operator to __half2 data type.
__host__ __device__ inline explicit operator
__nv_bfloat162() const
Conversion operator to __nv_bfloat162 data type.
__host__ __device__ inline explicit operator float2() const
Conversion operator to float2 data type.
Public Members
__nv_fp8x2_storage_t __x
Storage variable contains the vector of two scale factor values.
450
Chapter 15. Structs
15.24. __nv_fp8x4_e4m3
struct __nv_fp8x4_e4m3
__nv_fp8x4_e4m3 datatype
This structure implements the datatype for storage and operations on the vector of four fp8
values of e4m3 kind each: with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. The
encoding doesn’t support Infinity. NaNs are limited to 0x7F and 0xFF values.
Public Functions
__nv_fp8x4_e4m3() = default
Constructor by default.
__host__ __device__ inline explicit __nv_fp8x4_e4m3(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior
for out-of-range values.
__host__ __device__ inline explicit __nv_fp8x4_e4m3(const __nv_bfloat162 flo, const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
__host__ __device__ inline explicit __nv_fp8x4_e4m3(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp8x4_e4m3(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit operator float4() const
Conversion operator to float4 vector data type.
Public Members
__nv_fp8x4_storage_t __x
Storage variable contains the vector of four fp8 floating-point data values.
15.25. __nv_fp8x4_e5m2
struct __nv_fp8x4_e5m2
__nv_fp8x4_e5m2 datatype
This structure implements the datatype for handling four fp8 floating-point numbers of e5m2
kind each: with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits.
The structure implements converting constructors and operators.
15.24. __nv_fp8x4_e4m3
451
Public Functions
__nv_fp8x4_e5m2() = default
Constructor by default.
__host__ __device__ inline explicit __nv_fp8x4_e5m2(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior
for out-of-range values.
__host__ __device__ inline explicit __nv_fp8x4_e5m2(const __nv_bfloat162 flo, const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
__host__ __device__ inline explicit __nv_fp8x4_e5m2(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp8x4_e5m2(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit operator float4() const
Conversion operator to float4 vector data type.
Public Members
__nv_fp8x4_storage_t __x
Storage variable contains the vector of four fp8 floating-point data values.
15.26. __nv_fp8x4_e8m0
struct __nv_fp8x4_e8m0
__nv_fp8x4_e8m0 datatype
This structure implements the datatype for storage and operations on the vector of scale factors
of e8m0 kind each.
Public Functions
__nv_fp8x4_e8m0() = default
Constructor by default.
__host__ __device__ inline explicit __nv_fp8x4_e8m0(const __half2 flo, const __half2 fhi)
Constructor from a pair of __half2 data type values, relies on __NV_SATFINITE behavior
for out-of-range values.
452
Chapter 15. Structs
__host__ __device__ inline explicit __nv_fp8x4_e8m0(const __nv_bfloat162 flo, const
__nv_bfloat162 fhi)
Constructor from a pair of __nv_bfloat162 data type values, relies on __NV_SATFINITE
behavior for out-of-range values.
__host__ __device__ inline explicit __nv_fp8x4_e8m0(const double4 f)
Constructor from double4 vector data type, relies on __NV_SATFINITE behavior for out-
of-range values.
__host__ __device__ inline explicit __nv_fp8x4_e8m0(const float4 f)
Constructor from float4 vector data type, relies on __NV_SATFINITE behavior for out-of-
range values.
__host__ __device__ inline explicit operator float4() const
Conversion operator to float4 vector data type.
Public Members
__nv_fp8x4_storage_t __x
Storage variable contains the vector of four scale factor values.
__half
__half data type
__half2
__half2 data type
__half2_raw
__half2_raw data type
__half_raw
__half_raw data type
__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
__nv_fp4_e2m1
__nv_fp4_e2m1 datatype
__nv_fp4x2_e2m1
__nv_fp4x2_e2m1 datatype
__nv_fp4x4_e2m1
__nv_fp4x4_e2m1 datatype
__nv_fp6_e2m3
__nv_fp6_e2m3 datatype
__nv_fp6_e3m2
__nv_fp6_e3m2 datatype
15.26. __nv_fp8x4_e8m0
453
__nv_fp6x2_e2m3
__nv_fp6x2_e2m3 datatype
__nv_fp6x2_e3m2
__nv_fp6x2_e3m2 datatype
__nv_fp6x4_e2m3
__nv_fp6x4_e2m3 datatype
__nv_fp6x4_e3m2
__nv_fp6x4_e3m2 datatype
__nv_fp8_e4m3
__nv_fp8_e4m3 datatype
__nv_fp8_e5m2
__nv_fp8_e5m2 datatype
__nv_fp8_e8m0
__nv_fp8_e8m0 datatype
__nv_fp8x2_e4m3
__nv_fp8x2_e4m3 datatype
__nv_fp8x2_e5m2
__nv_fp8x2_e5m2 datatype
__nv_fp8x2_e8m0
__nv_fp8x2_e8m0 datatype
__nv_fp8x4_e4m3
__nv_fp8x4_e4m3 datatype
__nv_fp8x4_e5m2
__nv_fp8x4_e5m2 datatype
__nv_fp8x4_e8m0
__nv_fp8x4_e8m0 datatype
454
Chapter 15. Structs
Chapter 16. Notices
16.1. Notice
This document is provided for information purposes only and shall not be regarded as a warranty of a
certain functionality, condition, or quality of a product. NVIDIA Corporation (“NVIDIA”) makes no repre-
sentations or warranties, expressed or implied, as to the accuracy or completeness of the information
contained in this document and assumes no responsibility for any errors contained herein. NVIDIA shall
have no liability for the consequences or use of such information or for any infringement of patents
or other rights of third parties that may result from its use. This document is not a commitment to
develop, release, or deliver any Material (defined below), code, or functionality.
NVIDIA reserves the right to make corrections, modifications, enhancements, improvements, and any
other changes to this document, at any time without notice.
Customer should obtain the latest relevant information before placing orders and should verify that
such information is current and complete.
NVIDIA products are sold subject to the NVIDIA standard terms and conditions of sale supplied at the
time of order acknowledgement, unless otherwise agreed in an individual sales agreement signed by
authorized representatives of NVIDIA and customer (“Terms of Sale”). NVIDIA hereby expressly objects
to applying any customer general terms and conditions with regards to the purchase of the NVIDIA
product referenced in this document. No contractual obligations are formed either directly or indirectly
by this document.
NVIDIA products are not designed, authorized, or warranted to be suitable for use in medical, military,
aircraft, space, or life support equipment, nor in applications where failure or malfunction of the NVIDIA
product can reasonably be expected to result in personal injury, death, or property or environmental
damage. NVIDIA accepts no liability for inclusion and/or use of NVIDIA products in such equipment or
applications and therefore such inclusion and/or use is at customer’s own risk.
NVIDIA makes no representation or warranty that products based on this document will be suitable for
any specified use. Testing of all parameters of each product is not necessarily performed by NVIDIA.
It is customer’s sole responsibility to evaluate and determine the applicability of any information con-
tained in this document, ensure the product is suitable and fit for the application planned by customer,
and perform the necessary testing for the application in order to avoid a default of the application or
the product. Weaknesses in customer’s product designs may affect the quality and reliability of the
NVIDIA product and may result in additional or different conditions and/or requirements beyond those
contained in this document. NVIDIA accepts no liability related to any default, damage, costs, or prob-
lem which may be based on or attributable to: (i) the use of the NVIDIA product in any manner that is
contrary to this document or (ii) customer product designs.
No license, either expressed or implied, is granted under any NVIDIA patent right, copyright, or other
NVIDIA intellectual property right under this document. Information published by NVIDIA regarding
third-party products or services does not constitute a license from NVIDIA to use such products or
455
services or a warranty or endorsement thereof. Use of such information may require a license from a
third party under the patents or other intellectual property rights of the third party, or a license from
NVIDIA under the patents or other intellectual property rights of NVIDIA.
Reproduction of information in this document is permissible only if approved in advance by NVIDIA
in writing, reproduced without alteration and in full compliance with all applicable export laws and
regulations, and accompanied by all associated conditions, limitations, and notices.
THIS DOCUMENT AND ALL NVIDIA DESIGN SPECIFICATIONS, REFERENCE BOARDS, FILES, DRAWINGS,
DIAGNOSTICS, LISTS, AND OTHER DOCUMENTS (TOGETHER AND SEPARATELY, “MATERIALS”) ARE
BEING PROVIDED “AS IS.” NVIDIA MAKES NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR
OTHERWISE WITH RESPECT TO THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WAR-
RANTIES OF NONINFRINGEMENT, MERCHANTABILITY, AND FITNESS FOR A PARTICULAR PURPOSE.
TO THE EXTENT NOT PROHIBITED BY LAW, IN NO EVENT WILL NVIDIA BE LIABLE FOR ANY DAMAGES,
INCLUDING WITHOUT LIMITATION ANY DIRECT, INDIRECT, SPECIAL, INCIDENTAL, PUNITIVE, OR CON-
SEQUENTIAL DAMAGES, HOWEVER CAUSED AND REGARDLESS OF THE THEORY OF LIABILITY, ARIS-
ING OUT OF ANY USE OF THIS DOCUMENT, EVEN IF NVIDIA HAS BEEN ADVISED OF THE POSSIBILITY
OF SUCH DAMAGES. Notwithstanding any damages that customer might incur for any reason whatso-
ever, NVIDIA’s aggregate and cumulative liability towards customer for the products described herein
shall be limited in accordance with the Terms of Sale for the product.
16.2. OpenCL
OpenCL is a trademark of Apple Inc. used under license to the Khronos Group Inc.
16.3. Trademarks
NVIDIA and the NVIDIA logo are trademarks or registered trademarks of NVIDIA Corporation in the
U.S. and other countries. Other company and product names may be trademarks of the respective
companies with which they are associated.
©2007-2025, NVIDIA Corporation & affiliates. All rights reserved
456
Chapter 16. Notices
|