在 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 位。但是，如果两个线程想要更​​改数组第一个字符的不同位怎么办？即使位内的更改将位于不同的位置，他们也必须以原子方式执行操作。因此，使用原子操作可能会破坏并行性，在这种情况下，不需要使用原子操作，它没有任何意义，但必须使用，因为使用字符数组而不是更专门的东西的限制像一个std::bitset

先感谢您。

Answer 1

我对此问题提供了较晚的答案，以将其从未回答的列表中删除。

要实现您想要实现的目标，您可以定义一个unsigned int长度为 s 的数组N/32，其中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": http://stackoverflow.com/questions/9280654/c-printing-bits
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();
}