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) - v10.0.130 (older) - Last updated August 25, 2018 - Send Feedback

## 1.7. Integer Intrinsics

This section describes integer intrinsic functions that are only supported in device code. To use these functions you do not need to include any additional header files in your program.

### 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__ ​ unsigned int __funnelshift_l ( unsigned int lo, unsigned int hi, unsigned int shift )
  Concatenate `hi` : `lo`, shift left by `shift` & 31 bits, return the most significant 32 bits.
__device__ ​ unsigned int __funnelshift_lc ( unsigned int lo, unsigned int hi, unsigned int shift )
  Concatenate `hi` : `lo`, shift left by min(`shift`, 32) bits, return the most significant 32 bits.
__device__ ​ unsigned int __funnelshift_r ( unsigned int lo, unsigned int hi, unsigned int shift )
  Concatenate `hi` : `lo`, shift right by `shift` & 31 bits, return the least significant 32 bits.
__device__ ​ unsigned int __funnelshift_rc ( unsigned int lo, unsigned int hi, unsigned int shift )
  Concatenate `hi` : `lo`, shift right by min(`shift`, 32) bits, return the least significant 32 bits.
__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__ ​ unsigned int __funnelshift_l ( unsigned int lo, unsigned int hi, unsigned int shift )
  Concatenate `hi` : `lo`, shift left by `shift` & 31 bits, return the most significant 32 bits.

  ###### Returns

  Returns the most significant 32 bits of the shifted 64-bit value.

  ###### Description

  Shift the 64-bit value formed by concatenating argument `lo` and `hi` left by the amount specified by the argument `shift`. Argument `lo` holds bits 31:0 and argument `hi` holds bits 63:32 of the 64-bit source value. The source is shifted left by the wrapped value of `shift` (`shift` & 31). The most significant 32-bits of the result are returned.

__device__ ​ unsigned int __funnelshift_lc ( unsigned int lo, unsigned int hi, unsigned int shift )
  Concatenate `hi` : `lo`, shift left by min(`shift`, 32) bits, return the most significant 32 bits.

  ###### Returns

  Returns the most significant 32 bits of the shifted 64-bit value.

  ###### Description

  Shift the 64-bit value formed by concatenating argument `lo` and `hi` left by the amount specified by the argument `shift`. Argument `lo` holds bits 31:0 and argument `hi` holds bits 63:32 of the 64-bit source value. The source is shifted left by the clamped value of `shift` (min(`shift`, 32)). The most significant 32-bits of the result are returned.

__device__ ​ unsigned int __funnelshift_r ( unsigned int lo, unsigned int hi, unsigned int shift )
  Concatenate `hi` : `lo`, shift right by `shift` & 31 bits, return the least significant 32 bits.

  ###### Returns

  Returns the least significant 32 bits of the shifted 64-bit value.

  ###### Description

  Shift the 64-bit value formed by concatenating argument `lo` and `hi` right by the amount specified by the argument `shift`. Argument `lo` holds bits 31:0 and argument `hi` holds bits 63:32 of the 64-bit source value. The source is shifted right by the wrapped value of `shift` (`shift` & 31). The least significant 32-bits of the result are returned.

__device__ ​ unsigned int __funnelshift_rc ( unsigned int lo, unsigned int hi, unsigned int shift )
  Concatenate `hi` : `lo`, shift right by min(`shift`, 32) bits, return the least significant 32 bits.

  ###### Returns

  Returns the least significant 32 bits of the shifted 64-bit value.

  ###### Description

  Shift the 64-bit value formed by concatenating argument `lo` and `hi` right by the amount specified by the argument `shift`. Argument `lo` holds bits 31:0 and argument `hi` holds bits 63:32 of the 64-bit source value. The source is shifted right by the clamped value of `shift` (min(`shift`, 32)). The least significant 32-bits of the result are returned.

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