13. 整数内建函数

本节介绍整数内置函数。

所有这些函数在设备代码中都受支持。对于其中一些函数，还提供了特定于主机的实现。例如，参见__nv_bswap16()。要使用这些函数，您无需在程序中包含任何额外的头文件。

Functions

__device__ unsigned int __brev(unsigned int x)

反转一个32位无符号整数的比特位顺序。

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

反转64位无符号整数的位顺序。

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

从两个32位无符号整数中返回选定的字节。

__device__ int __clz(int x)

返回一个32位整数中连续高位零位的数量。

__device__ int __clzll(long long int x)

统计64位整数中连续高位零位的数量。

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

双向signed int16与int8的点积运算，使用int32累加器，并取第二个输入的高半部分。

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

双向unsigned int16与int8的点积运算，采用unsigned int32累加方式，并取第二个输入的上半部分。

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

双向unsigned int16与int8的点积运算，采用unsigned int32累加方式，并取第二个输入的高半部分。

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

双向signed int16与int8的点积运算，使用int32累加器，并取第二个输入的高半部分。

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

双向unsigned int16与int8的点积运算，使用unsigned int32累加，并取第二个输入的低半部分。

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

双向signed int16与int8的点积运算，使用int32累加，并取第二个输入的低半部分。

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

双向unsigned int16与int8的点积运算，使用unsigned int32累加器，并取第二个输入的低半部分。

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

双向signed int16与int8的点积运算，使用int32累加器，并取第二个输入的低半部分。

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

四路 unsigned int8 点积运算，使用 unsigned int32 累加器。

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

四路 unsigned int8 点积运算，使用 unsigned int32 累加器。

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

四路signed int8点积运算，使用int32累加器。

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

四路signed int8点积运算，使用int32累加器。

__device__ int __ffs(int x)

找出32位整数中最低有效位设置为1的位置。

__device__ int __ffsll(long long int x)

查找64位整数中最低有效位设置为1的位置。

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

在32位整数中查找第n个被设置为1的位的位置。

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

将hi与lo连接，左移shift & 31位，返回最高有效的32位。

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

将hi与lo连接，左移min(shift, 32)位，返回最高有效的32位。

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

将hi与lo连接，右移shift & 31位，返回最低有效的32位。

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

将hi与lo连接，然后右移min(shift, 32)位，返回最低有效的32位。

__device__ int __hadd(int x, int y)

计算有符号输入参数的平均值，避免中间求和时溢出。

__device__ int __mul24(int x, int y)

计算两个整数最低24位乘积的最低32位。

__device__ long long int __mul64hi(long long int x, long long int y)

计算两个64位整数乘积的最高有效64位。

__device__ int __mulhi(int x, int y)

计算两个32位整数乘积的最高有效32位。

__host__ __device__ unsigned short __nv_bswap16(unsigned short x)

反转16位无符号整数的字节顺序。

__host__ __device__ unsigned int __nv_bswap32(unsigned int x)

反转32位无符号整数的字节顺序。

__host__ __device__ unsigned long long __nv_bswap64(unsigned long long x)

反转64位无符号整数的字节顺序。

__device__ int __popc(unsigned int x)

计算32位整数中设置为1的位数。

__device__ int __popcll(unsigned long long int x)

计算一个64位整数中设置为1的位数。

__device__ int __rhadd(int x, int y)

计算有符号输入参数的舍入平均值，避免中间求和时发生溢出。

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

计算 \(|x - y| + z\) ，即绝对差之和。

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

计算无符号输入参数的平均值，避免中间求和时发生溢出。

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

计算两个无符号整数最低24位乘积的最低32位。

__device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y)

计算两个64位无符号整数乘积的最高有效64位。

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

计算两个32位无符号整数乘积的最高有效32位。

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

计算无符号输入参数的舍入平均值，避免中间求和时溢出。

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

计算 \(|x - y| + z\) ，即绝对差之和。

13.1. 函数

__device__ unsigned int __brev(unsigned int x)

反转一个32位无符号整数的比特位顺序。

反转32位无符号整数x的位顺序。

Returns

返回x的位反转值。即返回值的第N位对应x的第31-N位。

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

反转一个64位无符号整数的比特位顺序。

反转64位无符号整数x的位顺序。

Returns

返回x的位反转值。即返回值的第N位对应x的第63-N位。

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

从两个32位无符号整数中返回选定的字节。

创建8字节源数据

• uint64_t tmp64 = ((uint64_t)y << 32) | x;

提取选择器位

• selector0 = (s >> 0) & 0x7;

• selector1 = (s >> 4) & 0x7;

• selector2 = (s >> 8) & 0x7;

• selector3 = (s >> 12) & 0x7;

从8字节源数据中返回4个选定的字节：

• res[07:00] = tmp64[selector0];

• res[15:08] = tmp64[selector1];

• res[23:16] = tmp64[selector2];

• res[31:24] = tmp64[selector3];

Returns

根据选择器s的指定，从两个输入整数x和y提供的八个输入字节中返回一个由四个字节组成的32位整数。

__device__ int __clz(int x)

返回一个32位整数中连续高位零位的数量。

计算从x的最高有效位(第31位)开始的连续前导零位的数量。

Returns

返回一个介于0到32之间的值（包含0和32），表示零位的数量。

__device__ int __clzll(long long int x)

统计64位整数中连续高位零位的数量。

计算从x的最高有效位(第63位)开始的连续前导零位的数量。

Returns

返回一个介于0到64之间的值，表示零位的数量。

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

双向signed int16与int8的点积运算，采用int32累加器，并取第二个输入的高半部分。

从scrA中提取两个打包的16位整数，并从srcB的高16位中提取两个打包的8位整数，然后创建两个成对的8x16乘积，并将它们相加得到一个有符号的32位整数c。

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

双向unsigned int16与int8的点积运算，使用unsigned int32累加，并取第二个输入的上半部分。

从scrA中提取两个打包的16位整数，并从srcB的高16位中提取两个打包的8位整数，然后创建两个成对的8x16乘积并将它们相加，得到一个无符号32位整数c。

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

双向unsigned int16与int8的点积运算，采用unsigned int32累加，并取第二个输入的上半部分。

从scrA向量中取出两个打包的16位整数，并从srcB向量的高16位中取出两个打包的8位整数，然后创建两个成对的8x16乘积，并将它们相加为一个无符号32位整数c。

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

双向signed int16与int8的点积运算，使用int32累加，并取第二个输入的上半部分。

从scrA向量中取出两个打包的16位整数，并从srcB向量的高16位中取出两个打包的8位整数，然后创建两个成对的8x16乘积并将它们相加，得到一个有符号的32位整数c。

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

双向unsigned int16与int8的点积运算，采用unsigned int32累加方式，并取第二个输入的低半部分。

从scrA向量中取出两个打包的16位整数，并从srcB向量的低16位中取出两个打包的8位整数，然后创建两个成对的8x16乘积，并将它们相加为一个无符号32位整数c。

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

双向signed int16与int8的点积运算，使用int32累加器，并取第二个输入的低半部分。

从scrA向量中取出两个打包的16位整数，并从srcB向量的低16位中取出两个打包的8位整数，然后创建两个成对的8x16乘积并将它们相加，得到一个带符号的32位整数c。

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

双向unsigned int16与int8的点积运算，使用unsigned int32累加，并取第二个输入的低半部分。

从scrA中提取两个打包的16位整数，并从srcB的低16位中提取两个打包的8位整数，然后创建两个成对的8x16乘积并将它们相加，得到一个无符号32位整数c。

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

双向signed int16与int8的点积运算，使用int32累加器，并取第二个输入的低半部分。

从scrA中提取两个打包的16位整数，并从srcB的低16位中提取两个打包的8位整数，然后创建两个成对的8x16乘积并将它们相加，得到一个有符号的32位整数c。

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

四路 unsigned int8 点积运算，使用 unsigned int32 累加器。

从scrA和srcB向量中取出四对打包的字节大小整数，然后创建四个成对乘积并将它们相加，得到一个无符号32位整数c。

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

四路 unsigned int8 点积运算，使用 unsigned int32 累加器。

从scrA和srcB中提取四对打包的字节大小整数，然后创建四个成对乘积并将它们相加到一个无符号32位整数c中。

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

四路signed int8点积运算，使用int32累加器。

从scrA和srcB中提取四对打包的字节大小整数，然后创建四个成对乘积并将它们相加，得到一个有符号的32位整数c。

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

四路signed int8点积运算，使用int32累加器。

从scrA和srcB向量中取出四对打包的字节大小整数，然后创建四个成对乘积并将它们相加，得到一个有符号的32位整数c。

__device__ int __ffs(int x)

查找32位整数中最低有效位设置为1的位置。

查找x中第一个(最低有效位)设置为1的位的位置，其中最低有效位的位置为1。

Returns

返回一个介于0到32之间的值（包含0和32），表示第一个置位比特的位置。

• __ffs(0) 返回 0。

__device__ int __ffsll(long long int x)

查找64位整数中最低有效位设置为1的位置。

查找x中第一个(最低有效位)设置为1的位的位置，其中最低有效位的位置为1。

Returns

返回一个介于0到64之间的值（包含0和64），表示第一个置位比特的位置。

• __ffsll(0) 返回 0。

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

在一个32位整数中查找第n个被置为1的位的位置。

给定一个32位值mask和一个整数值base(范围0到31)，从base位开始查找mask中第n个(由offset指定)被置位的位。如果未找到，返回0xFFFFFFFF。

See also https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-fns for more information.

Returns

返回一个介于0到32之间的值（包含0和32），表示第n个置位的位置。

• 参数 base 必须小于等于31，否则行为未定义。

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

将hi与lo连接，左移shift & 31位，返回最高32位。

将由参数lo和hi拼接形成的64位值，按照参数shift指定的数量向左移位。参数lo包含64位源值的31:0位，参数hi包含63:32位。源值按照shift的环绕值(shift & 31)向左移位。返回结果的高32位。

Returns

返回移位后的64位值中最重要的32位。

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

将hi与lo连接，左移min(shift, 32)位，返回最高有效的32位。

将由参数lo和hi拼接形成的64位值，按照参数shift指定的数量向左移位。参数lo包含64位源值的31:0位，参数hi包含63:32位。源值按照shift的截断值（min(shift, 32)）向左移位。返回结果的高32位。

Returns

返回移位后的64位值中最重要的32位。

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

将hi与lo连接，右移shift & 31位，返回最低有效的32位。

将由参数lo和hi拼接形成的64位值，按照参数shift指定的数量向右移位。参数lo包含64位源值的31:0位，参数hi包含63:32位。源值将按照shift的包装值（shift & 31）向右移位。返回结果的低32位。

Returns

返回移位后的64位值的最低有效32位。

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

将hi与lo连接，右移min(shift, 32)位，返回最低有效的32位。

将由参数lo和hi拼接形成的64位值，按照参数shift指定的数量向右移位。参数lo包含64位源值的31:0位，参数hi包含63:32位。源值将按照shift的截断值（min(shift, 32)）向右移位。返回结果的低32位。

Returns

返回移位后的64位值的最低有效32位。

__device__ int __hadd(int x, int y)

计算有符号输入参数的平均值，避免中间求和时溢出。

计算有符号输入参数x和y的平均值为( x + y ) >> 1，避免中间求和时溢出。

Returns

返回一个有符号整数值，表示两个输入的有符号平均值。

__device__ int __mul24(int x, int y)

计算两个整数最低24位乘积的最低32位。

计算x和y最低有效24位乘积的最低有效32位。x和y的高8位将被忽略。

Returns

返回乘积 x * y 的最低有效32位。

__device__ long long int __mul64hi(long long int x, long long int y)

计算两个64位整数乘积的最高有效64位。

计算128位乘积x * y的最高有效64位，其中x和y是64位整数。

Returns

返回乘积 x * y 的最高有效64位。

__device__ int __mulhi(int x, int y)

计算两个32位整数乘积的最高有效32位。

计算64位乘积x * y的最高有效32位，其中x和y是32位整数。

Returns

返回乘积 x * y 的最高32位。

__host__ __device__ unsigned short __nv_bswap16(unsigned short x)

反转16位无符号整数的字节顺序。

反转x的字节顺序。仅支持在MSVC和其他定义了__GNUC__宏的主机编译器中使用，例如GCC和CLANG。

Returns

返回字节顺序反转后的 x。

__host__ __device__ unsigned int __nv_bswap32(unsigned int x)

反转32位无符号整数的字节顺序。

反转x的字节顺序。仅在MSVC和其他定义了__GNUC__宏的主机编译器（如GCC和CLANG）中支持。

Returns

返回字节顺序反转后的 x。

__host__ __device__ unsigned long long __nv_bswap64(unsigned long long x)

反转64位无符号整数的字节顺序。

反转x的字节顺序。仅支持在MSVC和其他定义了__GNUC__宏的主机编译器中使用，例如GCC和CLANG。

Returns

返回字节顺序反转后的 x。

__device__ int __popc(unsigned int x)

计算32位整数中设置为1的位数。

统计x中设置为1的比特位数。

Returns

返回一个介于0到32之间的值，表示设置的位数。

__device__ int __popcll(unsigned long long int x)

统计64位整数中设置为1的比特数量。

统计x中设置为1的比特位数。

Returns

返回一个介于0到64之间的值，表示设置的位数。

__device__ int __rhadd(int x, int y)

计算有符号输入参数的舍入平均值，避免中间求和过程中的溢出。

计算有符号输入参数x和y的平均值为( x + y + 1 ) >> 1，避免中间求和时发生溢出。

Returns

返回一个有符号整数值，表示两个输入的有符号四舍五入平均值。

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

计算 \( |x - y| + z \) ，即绝对差之和。

计算 \( |x - y| + z \)，即第三个参数 z 加上第一个参数 x 与第二个参数 y 差值的绝对值的32位总和。

输入 x 和 y 是32位有符号整数，输入 z 是32位无符号整数。

Returns

返回 \( |x - y| + z \)。

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

计算无符号输入参数的平均值，避免中间求和时溢出。

计算无符号输入参数x和y的平均值为( x + y ) >> 1，避免中间求和时溢出。

Returns

返回一个无符号整数值，表示两个输入的无符号平均值。

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

计算两个无符号整数最低24位乘积的最低32位。

计算x和y最低有效24位乘积的最低有效32位。x和y的高8位将被忽略。

Returns

返回乘积 x * y 的最低有效32位。

__device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y)

计算两个64位无符号整数乘积的最高有效64位。

计算128位乘积x * y的最高有效64位，其中x和y是64位无符号整数。

Returns

返回乘积 x * y 的最高有效64位。

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

计算两个32位无符号整数乘积的最高有效32位。

计算64位乘积x * y的最高有效32位，其中x和y是32位无符号整数。

Returns

返回乘积 x * y 的最高32位。

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

计算无符号输入参数的舍入平均值，避免中间求和时溢出。

计算无符号输入参数x和y的平均值为( x + y + 1 ) >> 1，避免中间求和时溢出。

Returns

返回一个无符号整数值，表示两个输入的无符号舍入平均值。

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

计算 \( |x - y| + z \) ，即绝对差之和。

计算 \( |x - y| + z \)，即第三个参数 z 加上第一个参数 x 与第二个参数 y 之差的绝对值的32位总和。

输入 x、y 和 z 是无符号32位整数。

Returns

返回 \( |x - y| + z \)。