## 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.