在 CUDA 中模拟 std::bitset

Question

我有一个提供给内核的输入数组。每个线程处理数组的一个值，并根据规则更改该值或根本不更改它。

我想很快找出输入内存中是否有任何变化，如果有，我想很快找到发生变化的位置（输入数组的索引）。

我想使用类似位数组的东西。位的总数将等于线程的总数。每个线程只会操作一个位，因此最初这些位将设置为假，如果一个线程更改相应的输入值，该位将变为真。

为了更清楚，假设我们有一个名为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

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

要实现您想要实现的目标，您可以定义一个长度为N/32 的unsigned ints 数组，其中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();
}