在CUDA中模拟std::bitset

Question

我有一个输入数组，它被传递给了一个内核。每个线程都处理数组的一个值，并根据规则更改该值或不更改。

之后我想很快找出输入内存中是否有任何更改，在这种情况下，我希望能够快速找到更改发生的位置（即输入数组的索引）。

我考虑使用类似于位数组的东西。位的总数将等于线程的总数。每个线程只会操作一个位，因此最初位设置为false，如果线程更改相应的输入值，则位将变为true。

为了使其更加清晰，假设我们有一个名为A的输入数组。

1 9 3 9 4 5

位数组应该如下所示：

0 0 0 0 0 0

因此，我们将有6个线程处理输入数组。假设最终的输入数组将是

1 9 3 9 2 5

因此，最终的二进制数组将是：

0 0 0 0 1 0

我不想使用一个 bool 数组，因为每个值将占用1字节内存，这相当多，因为我只想使用位操作。

是否有可能实现这样的功能？

我考虑创建一个char数组，其中每个值的二进制值都为8位。但是，如果两个线程想要更改数组的第一个字符的不同位，则它们必须以原子方式进行操作，即使位内部的更改将发生在不同的位置。因此，使用原子操作可能会破坏并行性，在这种情况下，使用原子操作是没有必要的，但由于使用 char 数组而不是像 std::bitset 这样的更专业的东西的限制，必须使用原子操作。

预先感谢您。

- ksm001

---

这个问题看起来很像你的问题：http://stackoverflow.com/questions/11042816/how-to-create-large-bit-array-in-cuda - BenC

谢谢，我已经阅读了问题和答案，但是它并没有提到如何使用位数组或者CUDA中是否有类似std::bitset的东西。对我来说，使用bool数组不是一个好主意，因为我不能在GPU上使用太多内存。 - ksm001

Answer 1

我提供一个晚一点的答案来解决这个问题，以将其从未回答的列表中移除。

为了实现你想要达到的目标，你可以定义一个长度为N/32的unsigned int数组，其中N是你要比较的数组的长度。然后，你可以使用atomicAdd函数根据两个数组元素是否相等来写入这样一个数组的每个位。

下面我提供一个简单的例子：

#include <iostream>

#include <thrust\device_vector.h>

__device__ unsigned int __ballot_non_atom(int predicate)
{
    if (predicate != 0) return (1 << (threadIdx.x % 32));
    else return 0;
}

__global__ void check_if_equal_elements(float* d_vec1_ptr, float* d_vec2_ptr, unsigned int* d_result, int Num_Warps_per_Block)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    const unsigned int warp_num = threadIdx.x >> 5;

    atomicAdd(&d_result[warp_num+blockIdx.x*Num_Warps_per_Block],__ballot_non_atom(!(d_vec1_ptr[tid] == d_vec2_ptr[tid])));
}

// --- Credit to "C printing bits": https://dev59.com/uWox5IYBdhLWcg3wVCxy
void printBits(unsigned int num){
    unsigned int size = sizeof(unsigned int);
    unsigned int maxPow = 1<<(size*8-1);
    int i=0;
    for(;i<size;++i){
        for(;i<size*8;++i){
            // print last bit and shift left.
            printf("%u ",num&maxPow ? 1 : 0);
            num = num<<1;
        }       
    }
}

void main(void)
{
    const int N = 64;

    thrust::device_vector<float> d_vec1(N,1.f);
    thrust::device_vector<float> d_vec2(N,1.f);

    d_vec2[3] = 3.f;
    d_vec2[7] = 4.f;

    unsigned int Num_Threads_per_Block      = 64;
    unsigned int Num_Blocks_per_Grid        = 1;
    unsigned int Num_Warps_per_Block        = Num_Threads_per_Block/32;
    unsigned int Num_Warps_per_Grid         = (Num_Threads_per_Block*Num_Blocks_per_Grid)/32;

    thrust::device_vector<unsigned int> d_result(Num_Warps_per_Grid,0);

    check_if_equal_elements<<<Num_Blocks_per_Grid,Num_Threads_per_Block>>>((float*)thrust::raw_pointer_cast(d_vec1.data()),
                                                                          (float*)thrust::raw_pointer_cast(d_vec2.data()),
                                                                          (unsigned int*)thrust::raw_pointer_cast(d_result.data()),
                                                                          Num_Warps_per_Block);

    unsigned int val = d_result[1];
    printBits(val);
    val = d_result[0];
    printBits(val);

    getchar();
} 

- Vitality