bit twiddle help:扩展位以跟随给定的位掩码

ncl*_*ent 2 algorithm cuda bit-manipulation

我对"扩展位"的快速方法感兴趣,可以定义如下:

  1. B为具有n位的二进制数,即{0,1} ^ n中的B \
  2. PB中所有1 /真位的位置,即1 << p[i] & B == 1和| P | = k
  3. 对于另一个给定数字,A\in {0,1} ^ k,令ApA给定B的位扩展形式,这样Ap[j] == A[j] << p[j].
  4. "比特扩展"的结果是Ap.

几个例子:

  • 给定B:00 1 0 111 0,A:0110,则Ap应为00 0 0 110 0
  • 给定B:1 00 1 1 00 1,A:1101,则Ap应为1 00 1 0 00 1

以下是一个简单的算法,但我不禁感到有一种更快/更容易的方法.

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;
}
Run Code Online (Sandbox Code Playgroud)

nju*_*ffa 5

问题似乎是要求对BMI2指令进行CUDA仿真,该指令PDEP采用源操作数a,并根据掩码的1位的位置存放其位b.对于当前发运的GPU上的相同或类似操作,没有硬件支持; 也就是说,包括Maxwell架构.

基于给出的两个例子,我假设掩模b通常是稀疏的,并且我们可以通过仅迭代1位来最小化工作b.这可能会导致GPU上出现分歧,但在不知道特定用例的情况下,性能的确切权衡是未知的.就目前而言,我假设b与分歧的负面影响相比,掩码中的稀疏性对性能的影响更大.

在下面的仿真代码中,我减少了可能"昂贵"的移位操作的使用,而主要依赖于简单的ALU指令.在各种GPU上,执行移位指令的吞吐量低于简单整数运算.我在代码的关键路径上保留了一个移位,以避免被算术单元限制执行.如果需要,1U << i可以通过添加来替换表达式:引入一个在循环之前m初始化的变量,并1在每次循环时加倍.

基本思想是b依次隔离每个1位掩码(从最低端开始),并将其与第i位的值隔离a,并将结果合并到扩展目标中.使用1位后b,我们将其从掩码中删除,并迭代直到掩码变为零.

为了避免将第i位移位a到位,我们简单地将其隔离,然后通过简单的否定将其值复制到所有更高有效位,利用整数的二进制补码表示.

/* 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;
}
Run Code Online (Sandbox Code Playgroud)

上面提到的没有任何移位操作的变体如下所示:

/* 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;
}
Run Code Online (Sandbox Code Playgroud)

在下面的评论中,@ Evgeny Kluev指出PDEP在国际象棋程序设计网站上的无移位仿真看起来可能比我上面两个实现中的任何一个都要 ; 这似乎值得一试.

  • 几乎相同的算法发布在 [chessprogramming](https://chessprogramming.wikispaces.com/BMI2) 上。但他们根本不使用移位指令。 (2认同)