Bit twiddle 帮助:扩展位以遵循给定的位掩码
Bit twiddle help: Expanding bits to follow a given bitmask
我对 "expanding bits," 的快速方法很感兴趣,它可以定义如下:
- 设B为n位的二进制数,即B\in { 0,1}^n
- 设P为B中所有1/真位的位置,即
1 << p[i] & B == 1
,|P|=k
- 对于另一个给定的数字,A \in {0,1}^k,设Ap 是给定 B 的 A 的位扩展形式,使得
Ap[j] == A[j] << p[j]
.
- "bit expansion" 的结果是 Ap。
举几个例子:
- 给定 B: 0010 1110, A :0110,那么Ap应该是00001100
- 给定 B: 1001 1001,A:1101,那么Ap应该是1001 0001
以下是一个简单的算法,但我不禁觉得有一种 faster/easier 方法可以做到这一点。
unsigned int expand_bits(unsigned int A, unsigned int B, int n) {
int k = popcount(B); // cuda function, but there are good methods for this
unsigned int Ap = 0;
int j = k-1;
// Starting at the most significant bit,
for (int i = n - 1; i >= 0; --i) {
Ap <<= 1;
// if B is 1, add the value at A[j] to Ap, decrement j.
if (B & (1 << i)) {
Ap += (A >> j--) & 1;
}
}
return Ap;
}
问题似乎是要求 BMI2 指令 PDEP
的 CUDA 仿真,它采用源操作数 a
,并根据 1 位的位置存储其位一个面具 b
。在当前出货的 GPU 上,没有硬件支持相同或类似的操作;也就是说,直到并包括 Maxwell 架构。
我假设,基于给出的两个示例,掩码 b
通常是稀疏的,并且我们可以通过仅迭代 b
的 1 位来最小化工作量。这可能会导致 GPU 上的不同分支,但在不了解特定用例的情况下,性能的确切权衡是未知的。现在,我假设掩码中稀疏性的利用 b
与发散的负面影响相比对性能有更强的积极影响。
在下面的仿真代码中,我减少了可能 "expensive" 移位操作的使用,而是主要依赖简单的 ALU 指令。在各种 GPU 上,移位指令的执行吞吐量低于简单的整数运算。我保留了一个单一的转变,关闭代码的关键路径,以避免执行受到算术单元的限制。如果需要,表达式 1U << i
可以用加法代替:引入一个变量 m
,在循环之前初始化为 1
,并在每次循环中加倍。
基本思想是依次隔离掩码 b
的每个 1 位(从最低有效端开始),并将其与 i[=45= 的值a
的第 ] 位,并将结果合并到扩展目标中。使用 b
中的 1 位后,我们将其从掩码中移除,并迭代直到掩码变为零。
为了避免将 a
的第 i 位移动到位,我们简单地将它隔离,然后通过简单的方法将其值复制到所有更重要的位取反,利用整数的补码表示。
/* Emulate PDEP: deposit the bits of 'a' (starting with the least significant
bit) at the positions indicated by the set bits of the mask stored in 'b'.
*/
__device__ unsigned int my_pdep (unsigned int a, unsigned int b)
{
unsigned int l, s, r = 0;
int i;
for (i = 0; b; i++) { // iterate over 1-bits in mask, until mask becomes 0
l = b & (0 - b); // extract mask's least significant 1-bit
b = b ^ l; // clear mask's least significant 1-bit
s = 0 - (a & (1U << i)); // spread i-th bit of 'a' to more signif. bits
r = r | (l & s); // deposit i-th bit of 'a' at position of mask's 1-bit
}
return r;
}
上面提到的没有任何移位操作的变体如下所示:
/* Emulate PDEP: deposit the bits of 'a' (starting with the least significant
bit) at the positions indicated by the set bits of the mask stored in 'b'.
*/
__device__ unsigned int my_pdep (unsigned int a, unsigned int b)
{
unsigned int l, s, r = 0, m = 1;
while (b) { // iterate over 1-bits in mask, until mask becomes 0
l = b & (0 - b); // extract mask's least significant 1-bit
b = b ^ l; // clear mask's least significant 1-bit
s = 0 - (a & m); // spread i-th bit of 'a' to more significant bits
r = r | (l & s); // deposit i-th bit of 'a' at position of mask's 1-bit
m = m + m; // mask for next bit of 'a'
}
return r;
}
在下面的评论中,@Evgeny Kluev 在 chessprogramming 网站上指出了一个无偏移 PDEP
仿真,看起来 可能比我的任何一个 更快以上两个实现;看来值得一试。
我对 "expanding bits," 的快速方法很感兴趣,它可以定义如下:
- 设B为n位的二进制数,即B\in { 0,1}^n
- 设P为B中所有1/真位的位置,即
1 << p[i] & B == 1
,|P|=k - 对于另一个给定的数字,A \in {0,1}^k,设Ap 是给定 B 的 A 的位扩展形式,使得
Ap[j] == A[j] << p[j]
. - "bit expansion" 的结果是 Ap。
举几个例子:
- 给定 B: 0010 1110, A :0110,那么Ap应该是00001100
- 给定 B: 1001 1001,A:1101,那么Ap应该是1001 0001
以下是一个简单的算法,但我不禁觉得有一种 faster/easier 方法可以做到这一点。
unsigned int expand_bits(unsigned int A, unsigned int B, int n) {
int k = popcount(B); // cuda function, but there are good methods for this
unsigned int Ap = 0;
int j = k-1;
// Starting at the most significant bit,
for (int i = n - 1; i >= 0; --i) {
Ap <<= 1;
// if B is 1, add the value at A[j] to Ap, decrement j.
if (B & (1 << i)) {
Ap += (A >> j--) & 1;
}
}
return Ap;
}
问题似乎是要求 BMI2 指令 PDEP
的 CUDA 仿真,它采用源操作数 a
,并根据 1 位的位置存储其位一个面具 b
。在当前出货的 GPU 上,没有硬件支持相同或类似的操作;也就是说,直到并包括 Maxwell 架构。
我假设,基于给出的两个示例,掩码 b
通常是稀疏的,并且我们可以通过仅迭代 b
的 1 位来最小化工作量。这可能会导致 GPU 上的不同分支,但在不了解特定用例的情况下,性能的确切权衡是未知的。现在,我假设掩码中稀疏性的利用 b
与发散的负面影响相比对性能有更强的积极影响。
在下面的仿真代码中,我减少了可能 "expensive" 移位操作的使用,而是主要依赖简单的 ALU 指令。在各种 GPU 上,移位指令的执行吞吐量低于简单的整数运算。我保留了一个单一的转变,关闭代码的关键路径,以避免执行受到算术单元的限制。如果需要,表达式 1U << i
可以用加法代替:引入一个变量 m
,在循环之前初始化为 1
,并在每次循环中加倍。
基本思想是依次隔离掩码 b
的每个 1 位(从最低有效端开始),并将其与 i[=45= 的值a
的第 ] 位,并将结果合并到扩展目标中。使用 b
中的 1 位后,我们将其从掩码中移除,并迭代直到掩码变为零。
为了避免将 a
的第 i 位移动到位,我们简单地将它隔离,然后通过简单的方法将其值复制到所有更重要的位取反,利用整数的补码表示。
/* Emulate PDEP: deposit the bits of 'a' (starting with the least significant
bit) at the positions indicated by the set bits of the mask stored in 'b'.
*/
__device__ unsigned int my_pdep (unsigned int a, unsigned int b)
{
unsigned int l, s, r = 0;
int i;
for (i = 0; b; i++) { // iterate over 1-bits in mask, until mask becomes 0
l = b & (0 - b); // extract mask's least significant 1-bit
b = b ^ l; // clear mask's least significant 1-bit
s = 0 - (a & (1U << i)); // spread i-th bit of 'a' to more signif. bits
r = r | (l & s); // deposit i-th bit of 'a' at position of mask's 1-bit
}
return r;
}
上面提到的没有任何移位操作的变体如下所示:
/* Emulate PDEP: deposit the bits of 'a' (starting with the least significant
bit) at the positions indicated by the set bits of the mask stored in 'b'.
*/
__device__ unsigned int my_pdep (unsigned int a, unsigned int b)
{
unsigned int l, s, r = 0, m = 1;
while (b) { // iterate over 1-bits in mask, until mask becomes 0
l = b & (0 - b); // extract mask's least significant 1-bit
b = b ^ l; // clear mask's least significant 1-bit
s = 0 - (a & m); // spread i-th bit of 'a' to more significant bits
r = r | (l & s); // deposit i-th bit of 'a' at position of mask's 1-bit
m = m + m; // mask for next bit of 'a'
}
return r;
}
在下面的评论中,@Evgeny Kluev 在 chessprogramming 网站上指出了一个无偏移 PDEP
仿真,看起来 可能比我的任何一个 更快以上两个实现;看来值得一试。