10. 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 32bit unsigned integer.
 __device__ unsigned long long int __brevll(unsigned long long int x)

Reverse the bit order of a 64bit unsigned integer.
 __device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s)

Return selected bytes from two 32bit unsigned integers.
 __device__ int __clz(int x)

Return the number of consecutive highorder zero bits in a 32bit integer.
 __device__ int __clzll(long long int x)

Count the number of consecutive highorder zero bits in a 64bit integer.
 __device__ int __dp2a_hi(int srcA, int srcB, int c)

Twoway
signed
int16
byint8
dot product withint32
accumulate, taking the upper half of the second input.  __device__ unsigned int __dp2a_hi(unsigned int srcA, unsigned int srcB, unsigned int c)

Twoway
unsigned
int16
byint8
dot product withunsigned
int32
accumulate, taking the upper half of the second input.  __device__ unsigned int __dp2a_hi(ushort2 srcA, uchar4 srcB, unsigned int c)

Twoway
unsigned
int16
byint8
dot product withunsigned
int32
accumulate, taking the upper half of the second input.  __device__ int __dp2a_hi(short2 srcA, char4 srcB, int c)

Twoway
signed
int16
byint8
dot product withint32
accumulate, taking the upper half of the second input.  __device__ unsigned int __dp2a_lo(ushort2 srcA, uchar4 srcB, unsigned int c)

Twoway
unsigned
int16
byint8
dot product withunsigned
int32
accumulate, taking the lower half of the second input.  __device__ int __dp2a_lo(short2 srcA, char4 srcB, int c)

Twoway
signed
int16
byint8
dot product withint32
accumulate, taking the lower half of the second input.  __device__ unsigned int __dp2a_lo(unsigned int srcA, unsigned int srcB, unsigned int c)

Twoway
unsigned
int16
byint8
dot product withunsigned
int32
accumulate, taking the lower half of the second input.  __device__ int __dp2a_lo(int srcA, int srcB, int c)

Twoway
signed
int16
byint8
dot product withint32
accumulate, taking the lower half of the second input.  __device__ unsigned int __dp4a(uchar4 srcA, uchar4 srcB, unsigned int c)

Fourway
unsigned
int8
dot product withunsigned
int32
accumulate.  __device__ unsigned int __dp4a(unsigned int srcA, unsigned int srcB, unsigned int c)

Fourway
unsigned
int8
dot product withunsigned
int32
accumulate.  __device__ int __dp4a(int srcA, int srcB, int c)

Fourway
signed
int8
dot product withint32
accumulate.  __device__ int __dp4a(char4 srcA, char4 srcB, int c)

Fourway
signed
int8
dot product withint32
accumulate.  __device__ int __ffs(int x)

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

Find the position of the least significant bit set to 1 in a 64bit integer.
 __device__ unsigned __fns(unsigned mask, unsigned base, int offset)

Find the position of the nth set to 1 bit in a 32bit integer.
 __device__ unsigned int __funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift)

Concatenate
hi
:lo
, shift left byshift
& 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 byshift
& 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 x, int y)

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 64bit integers.
 __device__ int __mulhi(int x, int y)

Calculate the most significant 32 bits of the product of the two 32bit integers.
 __device__ int __popc(unsigned int x)

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

Count the number of bits that are set to 1 in a 64bit integer.
 __device__ int __rhadd(int x, int y)

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 x, unsigned int y)

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 32bit unsigned integers.
 __device__ unsigned int __urhadd(unsigned int x, unsigned int y)

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.
10.1. Functions

__device__ unsigned int __brev(unsigned int x)

Reverse the bit order of a 32bit unsigned integer.
Reverses the bit order of the 32bit unsigned integer
x
. Returns

Returns the bitreversed value of
x
. i.e. bit N of the return value corresponds to bit 31N ofx
.

__device__ unsigned long long int __brevll(unsigned long long int x)

Reverse the bit order of a 64bit unsigned integer.
Reverses the bit order of the 64bit unsigned integer
x
. Returns

Returns the bitreversed value of
x
. i.e. bit N of the return value corresponds to bit 63N ofx
.

__device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s)

Return selected bytes from two 32bit unsigned integers.
Create 8byte source
uint64_t
tmp64
= ((uint64_t)y
<< 32) x
;
Extract selector bits
selector0
= (s
>> 0) & 0x7;selector1
= (s
>> 4) & 0x7;selector2
= (s
>> 8) & 0x7;selector3
= (s
>> 12) & 0x7;
Return 4 selected bytes from 8byte source:
res
[07:00] =tmp64
[selector0
];res
[15:08] =tmp64
[selector1
];res
[23:16] =tmp64
[selector2
];res
[31:24] =tmp64
[selector3
];
 Returns

Returns a 32bit integer consisting of four bytes from eight input bytes provided in the two input integers
x
andy
, as specified by a selector,s
.

__device__ int __clz(int x)

Return the number of consecutive highorder zero bits in a 32bit integer.
Count the number of consecutive leading zero bits, starting at the most significant bit (bit 31) of
x
. Returns

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

__device__ int __clzll(long long int x)

Count the number of consecutive highorder zero bits in a 64bit integer.
Count the number of consecutive leading zero bits, starting at the most significant bit (bit 63) of
x
. Returns

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

__device__ int __dp2a_hi(int srcA, int srcB, int c)

Twoway
signed
int16
byint8
dot product withint32
accumulate, taking the upper half of the second input.Extracts two packed 16bit integers from
scrA
and two packed 8bit integers from the upper 16 bits ofsrcB
, then creates two pairwise 8x16 products and adds them together to a signed 32bit integerc
.

__device__ unsigned int __dp2a_hi(unsigned int srcA, unsigned int srcB, unsigned int c)

Twoway
unsigned
int16
byint8
dot product withunsigned
int32
accumulate, taking the upper half of the second input.Extracts two packed 16bit integers from
scrA
and two packed 8bit integers from the upper 16 bits ofsrcB
, then creates two pairwise 8x16 products and adds them together to an unsigned 32bit integerc
.

__device__ unsigned int __dp2a_hi(ushort2 srcA, uchar4 srcB, unsigned int c)

Twoway
unsigned
int16
byint8
dot product withunsigned
int32
accumulate, taking the upper half of the second input.Takes two packed 16bit integers from
scrA
vector and two packed 8bit integers from the upper 16 bits ofsrcB
vector, then creates two pairwise 8x16 products and adds them together to an unsigned 32bit integerc
.

__device__ int __dp2a_hi(short2 srcA, char4 srcB, int c)

Twoway
signed
int16
byint8
dot product withint32
accumulate, taking the upper half of the second input.Takes two packed 16bit integers from
scrA
vector and two packed 8bit integers from the upper 16 bits ofsrcB
vector, then creates two pairwise 8x16 products and adds them together to a signed 32bit integerc
.

__device__ unsigned int __dp2a_lo(ushort2 srcA, uchar4 srcB, unsigned int c)

Twoway
unsigned
int16
byint8
dot product withunsigned
int32
accumulate, taking the lower half of the second input.Takes two packed 16bit integers from
scrA
vector and two packed 8bit integers from the lower 16 bits ofsrcB
vector, then creates two pairwise 8x16 products and adds them together to an unsigned 32bit integerc
.

__device__ int __dp2a_lo(short2 srcA, char4 srcB, int c)

Twoway
signed
int16
byint8
dot product withint32
accumulate, taking the lower half of the second input.Takes two packed 16bit integers from
scrA
vector and two packed 8bit integers from the lower 16 bits ofsrcB
vector, then creates two pairwise 8x16 products and adds them together to a signed 32bit integerc
.

__device__ unsigned int __dp2a_lo(unsigned int srcA, unsigned int srcB, unsigned int c)

Twoway
unsigned
int16
byint8
dot product withunsigned
int32
accumulate, taking the lower half of the second input.Extracts two packed 16bit integers from
scrA
and two packed 8bit integers from the lower 16 bits ofsrcB
, then creates two pairwise 8x16 products and adds them together to an unsigned 32bit integerc
.

__device__ int __dp2a_lo(int srcA, int srcB, int c)

Twoway
signed
int16
byint8
dot product withint32
accumulate, taking the lower half of the second input.Extracts two packed 16bit integers from
scrA
and two packed 8bit integers from the lower 16 bits ofsrcB
, then creates two pairwise 8x16 products and adds them together to a signed 32bit integerc
.

__device__ unsigned int __dp4a(uchar4 srcA, uchar4 srcB, unsigned int c)

Fourway
unsigned
int8
dot product withunsigned
int32
accumulate.Takes four pairs of packed bytesized integers from
scrA
andsrcB
vectors, then creates four pairwise products and adds them together to an unsigned 32bit integerc
.

__device__ unsigned int __dp4a(unsigned int srcA, unsigned int srcB, unsigned int c)

Fourway
unsigned
int8
dot product withunsigned
int32
accumulate.Extracts four pairs of packed bytesized integers from
scrA
andsrcB
, then creates four pairwise products and adds them together to an unsigned 32bit integerc
.

__device__ int __dp4a(int srcA, int srcB, int c)

Fourway
signed
int8
dot product withint32
accumulate.Extracts four pairs of packed bytesized integers from
scrA
andsrcB
, then creates four pairwise products and adds them together to a signed 32bit integerc
.

__device__ int __dp4a(char4 srcA, char4 srcB, int c)

Fourway
signed
int8
dot product withint32
accumulate.Takes four pairs of packed bytesized integers from
scrA
andsrcB
vectors, then creates four pairwise products and adds them together to a signed 32bit integerc
.

__device__ int __ffs(int x)

Find the position of the least significant bit set to 1 in a 32bit integer.
Find the position of the first (least significant) bit set to 1 in
x
, where the least significant bit position is 1. Returns

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

__device__ int __ffsll(long long int x)

Find the position of the least significant bit set to 1 in a 64bit integer.
Find the position of the first (least significant) bit set to 1 in
x
, where the least significant bit position is 1. Returns

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

__device__ unsigned __fns(unsigned mask, unsigned base, int offset)

Find the position of the nth set to 1 bit in a 32bit integer.
Given a 32bit value
mask
and an integer valuebase
(between 0 and 31), find the nth (given byoffset
) set bit inmask
from thebase
bit. If not found, return 0xFFFFFFFF.See also https://docs.nvidia.com/cuda/parallelthreadexecution/index.html#integerarithmeticinstructionsfns for more information.
 Returns

Returns a value between 0 and 32 inclusive representing the position of the nth set bit.
parameter
base
must be <=31, otherwise behavior is undefined.

__device__ unsigned int __funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift)

Concatenate
hi
:lo
, shift left byshift
& 31 bits, return the most significant 32 bits.Shift the 64bit value formed by concatenating argument
lo
andhi
left by the amount specified by the argumentshift
. Argumentlo
holds bits 31:0 and argumenthi
holds bits 63:32 of the 64bit source value. The source is shifted left by the wrapped value ofshift
(shift
& 31). The most significant 32bits of the result are returned. Returns

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

__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.Shift the 64bit value formed by concatenating argument
lo
andhi
left by the amount specified by the argumentshift
. Argumentlo
holds bits 31:0 and argumenthi
holds bits 63:32 of the 64bit source value. The source is shifted left by the clamped value ofshift
(min(shift
, 32)). The most significant 32bits of the result are returned. Returns

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

__device__ unsigned int __funnelshift_r(unsigned int lo, unsigned int hi, unsigned int shift)

Concatenate
hi
:lo
, shift right byshift
& 31 bits, return the least significant 32 bits.Shift the 64bit value formed by concatenating argument
lo
andhi
right by the amount specified by the argumentshift
. Argumentlo
holds bits 31:0 and argumenthi
holds bits 63:32 of the 64bit source value. The source is shifted right by the wrapped value ofshift
(shift
& 31). The least significant 32bits of the result are returned. Returns

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

__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.Shift the 64bit value formed by concatenating argument
lo
andhi
right by the amount specified by the argumentshift
. Argumentlo
holds bits 31:0 and argumenthi
holds bits 63:32 of the 64bit source value. The source is shifted right by the clamped value ofshift
(min(shift
, 32)). The least significant 32bits of the result are returned. Returns

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

__device__ int __hadd(int x, int y)

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

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

__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.
Calculate the least significant 32 bits of the product of the least significant 24 bits of
x
andy
. The high order 8 bits ofx
andy
are ignored. Returns

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

__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 64bit integers.
Calculate the most significant 64 bits of the 128bit product
x
*y
, wherex
andy
are 64bit integers. Returns

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

__device__ int __mulhi(int x, int y)

Calculate the most significant 32 bits of the product of the two 32bit integers.
Calculate the most significant 32 bits of the 64bit product
x
*y
, wherex
andy
are 32bit integers. Returns

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

__device__ int __popc(unsigned int x)

Count the number of bits that are set to 1 in a 32bit integer.
Count the number of bits that are set to 1 in
x
. Returns

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

__device__ int __popcll(unsigned long long int x)

Count the number of bits that are set to 1 in a 64bit integer.
Count the number of bits that are set to 1 in
x
. Returns

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

__device__ int __rhadd(int x, int y)

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

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

__device__ unsigned int __sad(int x, int y, unsigned int z)

Calculate \( x  y + z \) , the sum of absolute difference.
Calculate \( x  y + z \) , the 32bit 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
andy
are signed 32bit integers, inputz
is a 32bit unsigned integer. Returns

Returns \( x  y + z \).

__device__ unsigned int __uhadd(unsigned int x, unsigned int y)

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

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

__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.
Calculate the least significant 32 bits of the product of the least significant 24 bits of
x
andy
. The high order 8 bits ofx
andy
are ignored. Returns

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

__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.
Calculate the most significant 64 bits of the 128bit product
x
*y
, wherex
andy
are 64bit unsigned integers. Returns

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

__device__ unsigned int __umulhi(unsigned int x, unsigned int y)

Calculate the most significant 32 bits of the product of the two 32bit unsigned integers.
Calculate the most significant 32 bits of the 64bit product
x
*y
, wherex
andy
are 32bit unsigned integers. Returns

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

__device__ unsigned int __urhadd(unsigned int x, unsigned int y)

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

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

__device__ unsigned int __usad(unsigned int x, unsigned int y, unsigned int z)

Calculate \( x  y + z \) , the sum of absolute difference.
Calculate \( x  y + z \) , the 32bit 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
, andz
are unsigned 32bit integers. Returns

Returns \( x  y + z \).