13. 整数内函数

本节介绍整数内联函数。

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

函数

__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 的位顺序。

返回值

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

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

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

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

返回值

返回 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];

返回值

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

__device__ int __clz(int x)

返回 32 位整数中前导零位的数量。

计算 x 的最高有效位（位 31）开始的连续前导零位的数量。

返回值

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

__device__ int __clzll(long long int x)

计算 64 位整数中前导零位的数量。

计算 x 的最高有效位（位 63）开始的连续前导零位的数量。

返回值

返回值介于 0 到 64 之间（包括 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。

返回值

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

• __ffs(0) 返回 0。

__device__ int __ffsll(long long int x)

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

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

返回值

返回值介于 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。

另请参阅 https://docs.nvda.net.cn/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-fns 以获取更多信息。

返回值

返回值介于 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 保存位 31:0，参数 hi 保存 64 位源值的位 63:32。源值左移 shift 的wrap值（shift & 31）。返回结果的最高有效 32 位。

返回值

返回移位后的 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 保存位 31:0，参数 hi 保存 64 位源值的位 63:32。源值左移 shift 的clamp值（min(shift, 32)）。返回结果的最高有效 32 位。

返回值

返回移位后的 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 保存位 31:0，参数 hi 保存 64 位源值的位 63:32。源值右移 shift 的wrap值（shift & 31）。返回结果的最低有效 32 位。

返回值

返回移位后的 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 保存位 31:0，参数 hi 保存 64 位源值的位 63:32。源值右移 shift 的clamp值（min(shift, 32)）。返回结果的最低有效 32 位。

返回值

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

__device__ int __hadd(int x, int y)

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

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

返回值

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

__device__ int __mul24(int x, int y)

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

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

返回值

返回乘积 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 位整数。

返回值

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

__device__ int __mulhi(int x, int y)

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

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

返回值

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

__host__ __device__ unsigned short __nv_bswap16(unsigned short x)

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

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

返回值

返回字节顺序反转的 x。

__host__ __device__ unsigned int __nv_bswap32(unsigned int x)

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

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

返回值

返回字节顺序反转的 x。

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

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

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

返回值

返回字节顺序反转的 x。

__device__ int __popc(unsigned int x)

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

计算 x 中设置为 1 的位数。

返回值

返回值介于 0 和 32 之间（包含 0 和 32），表示设置位的数量。

__device__ int __popcll(unsigned long long int x)

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

计算 x 中设置为 1 的位数。

返回值

返回值介于 0 和 64 之间（包含 0 和 64），表示设置位的数量。

__device__ int __rhadd(int x, int y)

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

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

返回值

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

__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 位无符号整数。

返回值

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

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

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

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

返回值

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

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

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

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

返回值

返回乘积 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 位无符号整数。

返回值

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

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

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

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

返回值

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

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

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

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

返回值

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

__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 位整数。

返回值

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