Эмуляция std :: bitset в CUDA

Question

У меня есть входной массив, который передается ядру. Каждый поток работает с одним значением массива и либо изменяет значение, либо не изменяет его вообще в соответствии с правилом.

Я бы хотел узнать об этом очень быстро, если бы произошли какие-либо изменения во входной памяти, и, если бы это было так, я бы очень быстро нашел, где произошло это изменение (индекс входного массива).

Я думал об использовании чего-то вроде массива бит. Общее количество бит будет равно общему количеству потоков. Каждый поток будет манипулировать только одним битом, поэтому изначально биты будут установлены на false, если поток изменит соответствующее входное значение, бит станет истинным.

Для того, чтобы сделать его более ясным, давайте предположим, что у нас есть этот входной массив называется 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 с длиной 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(); 
}