This document includes math equations (highlighted in red) which are best viewed with Firefox version 4.0 or higher, or another MathML-aware browser. There is also a PDF version of this document.
CUDA Math API (PDF) - v9.0.176 (older) - Last updated September 2, 2017 - Send Feedback

## 1.7. Integer Intrinsics

This section describes integer intrinsic functions that are only supported in device code.

### Functions

__device__ ​ unsigned int __brev ( unsigned int x )
  Reverse the bit order of a 32 bit unsigned integer.
__device__ ​ unsigned long long int __brevll ( unsigned long long int x )
  Reverse the bit order of a 64 bit unsigned integer.
__device__ ​ unsigned int __byte_perm ( unsigned int x, unsigned int y, unsigned int s )
  Return selected bytes from two 32 bit unsigned integers.
__device__ ​ int __clz ( int x )
  Return the number of consecutive high-order zero bits in a 32 bit integer.
__device__ ​ int __clzll ( long long int x )
  Count the number of consecutive high-order zero bits in a 64 bit integer.
__device__ ​ int __ffs ( int x )
  Find the position of the least significant bit set to 1 in a 32 bit integer.
__device__ ​ int __ffsll ( long long int x )
  Find the position of the least significant bit set to 1 in a 64 bit integer.
__device__ ​ int __hadd ( int , int )
  Compute average of signed input arguments, avoiding overflow in the intermediate sum.
__device__ ​ int __mul24 ( int x, int y )
  Calculate the least significant 32 bits of the product of the least significant 24 bits of two integers.
__device__ ​ long long int __mul64hi ( long long int x, long long int y )
  Calculate the most significant 64 bits of the product of the two 64 bit integers.
__device__ ​ int __mulhi ( int x, int y )
  Calculate the most significant 32 bits of the product of the two 32 bit integers.
__device__ ​ int __popc ( unsigned int x )
  Count the number of bits that are set to 1 in a 32 bit integer.
__device__ ​ int __popcll ( unsigned long long int x )
  Count the number of bits that are set to 1 in a 64 bit integer.
__device__ ​ int __rhadd ( int , int )
  Compute rounded average of signed input arguments, avoiding overflow in the intermediate sum.
__device__ ​ unsigned int __sad ( int x, int y, unsigned int z )
  Calculate | x − y | + z , the sum of absolute difference.
__device__ ​ unsigned int __uhadd ( unsigned int, unsigned int )
  Compute average of unsigned input arguments, avoiding overflow in the intermediate sum.
__device__ ​ unsigned int __umul24 ( unsigned int x, unsigned int y )
  Calculate the least significant 32 bits of the product of the least significant 24 bits of two unsigned integers.
__device__ ​ unsigned long long int __umul64hi ( unsigned long long int x, unsigned long long int y )
  Calculate the most significant 64 bits of the product of the two 64 unsigned bit integers.
__device__ ​ unsigned int __umulhi ( unsigned int x, unsigned int y )
  Calculate the most significant 32 bits of the product of the two 32 bit unsigned integers.
__device__ ​ unsigned int __urhadd ( unsigned int, unsigned int )
  Compute rounded average of unsigned input arguments, avoiding overflow in the intermediate sum.
__device__ ​ unsigned int __usad ( unsigned int x, unsigned int y, unsigned int z )
  Calculate | x − y | + z , the sum of absolute difference.

### Functions

__device__ ​ unsigned int __brev ( unsigned int x )
  Reverse the bit order of a 32 bit unsigned integer.

  ###### Returns

  Returns the bit-reversed value of `x`. i.e. bit N of the return value corresponds to bit 31-N of `x`.

  ###### Description

  Reverses the bit order of the 32 bit unsigned integer `x`.

__device__ ​ unsigned long long int __brevll ( unsigned long long int x )
  Reverse the bit order of a 64 bit unsigned integer.

  ###### Returns

  Returns the bit-reversed value of `x`. i.e. bit N of the return value corresponds to bit 63-N of `x`.

  ###### Description

  Reverses the bit order of the 64 bit unsigned integer `x`.

__device__ ​ unsigned int __byte_perm ( unsigned int x, unsigned int y, unsigned int s )
  Return selected bytes from two 32 bit unsigned integers.

  ###### Returns

  The returned value r is computed to be: `result[n] := input[selector[n]]` where `result[n]` is the nth byte of r.

  ###### Description

  byte_perm(x,y,s) returns a 32-bit integer consisting of four bytes from eight input bytes provided in the two input integers `x` and `y`, as specified by a selector, `s`.

  The input bytes are indexed as follows: input[0] = x<7:0> input[1] = x<15:8> input[2] = x<23:16> input[3] = x<31:24> input[4] = y<7:0> input[5] = y<15:8> input[6] = y<23:16> input[7] = y<31:24> The selector indices are as follows (the upper 16-bits of the selector are not used): selector[0] = s<2:0> selector[1] = s<6:4> selector[2] = s<10:8> selector[3] = s<14:12>

__device__ ​ int __clz ( int x )
  Return the number of consecutive high-order zero bits in a 32 bit integer.

  ###### Returns

  Returns a value between 0 and 32 inclusive representing the number of zero bits.

  ###### Description

  Count the number of consecutive leading zero bits, starting at the most significant bit (bit 31) of `x`.

__device__ ​ int __clzll ( long long int x )
  Count the number of consecutive high-order zero bits in a 64 bit integer.

  ###### Returns

  Returns a value between 0 and 64 inclusive representing the number of zero bits.

  ###### Description

  Count the number of consecutive leading zero bits, starting at the most significant bit (bit 63) of `x`.

__device__ ​ int __ffs ( int x )
  Find the position of the least significant bit set to 1 in a 32 bit integer.

  ###### Returns

  Returns a value between 0 and 32 inclusive representing the position of the first bit set.

  - __ffs(0) returns 0.

  ###### Description

  Find the position of the first (least significant) bit set to 1 in `x`, where the least significant bit position is 1.

__device__ ​ int __ffsll ( long long int x )
  Find the position of the least significant bit set to 1 in a 64 bit integer.

  ###### Returns

  Returns a value between 0 and 64 inclusive representing the position of the first bit set.

  - __ffsll(0) returns 0.

  ###### Description

  Find the position of the first (least significant) bit set to 1 in `x`, where the least significant bit position is 1.

__device__ ​ int __hadd ( int , int )
  Compute average of signed input arguments, avoiding overflow in the intermediate sum.

  ###### Returns

  Returns a signed integer value representing the signed average value of the two inputs.

  ###### Description

  Compute average of signed input arguments `x` and `y` as ( `x` + `y` ) >> 1, avoiding overflow in the intermediate sum.

__device__ ​ int __mul24 ( int x, int y )
  Calculate the least significant 32 bits of the product of the least significant 24 bits of two integers.

  ###### Returns

  Returns the least significant 32 bits of the product `x` * `y`.

  ###### Description

  Calculate the least significant 32 bits of the product of the least significant 24 bits of `x` and `y`. The high order 8 bits of `x` and `y` are ignored.

__device__ ​ long long int __mul64hi ( long long int x, long long int y )
  Calculate the most significant 64 bits of the product of the two 64 bit integers.

  ###### Returns

  Returns the most significant 64 bits of the product `x` * `y`.

  ###### Description

  Calculate the most significant 64 bits of the 128-bit product `x` * `y`, where `x` and `y` are 64-bit integers.

__device__ ​ int __mulhi ( int x, int y )
  Calculate the most significant 32 bits of the product of the two 32 bit integers.

  ###### Returns

  Returns the most significant 32 bits of the product `x` * `y`.

  ###### Description

  Calculate the most significant 32 bits of the 64-bit product `x` * `y`, where `x` and `y` are 32-bit integers.

__device__ ​ int __popc ( unsigned int x )
  Count the number of bits that are set to 1 in a 32 bit integer.

  ###### Returns

  Returns a value between 0 and 32 inclusive representing the number of set bits.

  ###### Description

  Count the number of bits that are set to 1 in `x`.

__device__ ​ int __popcll ( unsigned long long int x )
  Count the number of bits that are set to 1 in a 64 bit integer.

  ###### Returns

  Returns a value between 0 and 64 inclusive representing the number of set bits.

  ###### Description

  Count the number of bits that are set to 1 in `x`.

__device__ ​ int __rhadd ( int , int )
  Compute rounded average of signed input arguments, avoiding overflow in the intermediate sum.

  ###### Returns

  Returns a signed integer value representing the signed rounded average value of the two inputs.

  ###### Description

  Compute average of signed input arguments `x` and `y` as ( `x` + `y` + 1 ) >> 1, avoiding overflow in the intermediate sum.

__device__ ​ unsigned int __sad ( int x, int y, unsigned int z )
  Calculate | x − y | + z , the sum of absolute difference.

  ###### Returns

  Returns | x − y | + z .

  ###### Description

  Calculate | x − y | + z , the 32-bit sum of the third argument `z` plus and the absolute value of the difference between the first argument, `x`, and second argument, `y`.

  Inputs `x` and `y` are signed 32-bit integers, input `z` is a 32-bit unsigned integer.

__device__ ​ unsigned int __uhadd ( unsigned int, unsigned int )
  Compute average of unsigned input arguments, avoiding overflow in the intermediate sum.

  ###### Returns

  Returns an unsigned integer value representing the unsigned average value of the two inputs.

  ###### Description

  Compute average of unsigned input arguments `x` and `y` as ( `x` + `y` ) >> 1, avoiding overflow in the intermediate sum.

__device__ ​ unsigned int __umul24 ( unsigned int x, unsigned int y )
  Calculate the least significant 32 bits of the product of the least significant 24 bits of two unsigned integers.

  ###### Returns

  Returns the least significant 32 bits of the product `x` * `y`.

  ###### Description

  Calculate the least significant 32 bits of the product of the least significant 24 bits of `x` and `y`. The high order 8 bits of `x` and `y` are ignored.

__device__ ​ unsigned long long int __umul64hi ( unsigned long long int x, unsigned long long int y )
  Calculate the most significant 64 bits of the product of the two 64 unsigned bit integers.

  ###### Returns

  Returns the most significant 64 bits of the product `x` * `y`.

  ###### Description

  Calculate the most significant 64 bits of the 128-bit product `x` * `y`, where `x` and `y` are 64-bit unsigned integers.

__device__ ​ unsigned int __umulhi ( unsigned int x, unsigned int y )
  Calculate the most significant 32 bits of the product of the two 32 bit unsigned integers.

  ###### Returns

  Returns the most significant 32 bits of the product `x` * `y`.

  ###### Description

  Calculate the most significant 32 bits of the 64-bit product `x` * `y`, where `x` and `y` are 32-bit unsigned integers.

__device__ ​ unsigned int __urhadd ( unsigned int, unsigned int )
  Compute rounded average of unsigned input arguments, avoiding overflow in the intermediate sum.

  ###### Returns

  Returns an unsigned integer value representing the unsigned rounded average value of the two inputs.

  ###### Description

  Compute average of unsigned input arguments `x` and `y` as ( `x` + `y` + 1 ) >> 1, avoiding overflow in the intermediate sum.

__device__ ​ unsigned int __usad ( unsigned int x, unsigned int y, unsigned int z )
  Calculate | x − y | + z , the sum of absolute difference.

  ###### Returns

  Returns | x − y | + z .

  ###### Description

  Calculate | x − y | + z , the 32-bit sum of the third argument `z` plus and the absolute value of the difference between the first argument, `x`, and second argument, `y`.

  Inputs `x`, `y`, and `z` are unsigned 32-bit integers.