Emulation de std :: bitset dans CUDA

Question

J'ai un tableau d'entrée qui est donné à un noyau. Chaque thread fonctionne avec une valeur du tableau et modifie la valeur ou ne la modifie pas du tout selon une règle.

Je voudrais savoir très rapidement après s'il y avait des changements dans la mémoire d'entrée et, dans le cas où, je voudrais trouver très rapidement où ce changement s'est produit (index du tableau d'entrée). J'ai pensé à utiliser quelque chose comme un tableau de bits. La quantité totale de bits serait égale à la quantité totale de threads. Chaque thread ne manipulerait qu'un seul bit, donc initialement les bits seraient mis à false, si un thread change la valeur d'entrée correspondante, le bit deviendra vrai.

Afin de le rendre plus clair, supposons que nous avons ce tableau d'entrée appelé A

1 9 3 9 4 5 

Le tableau de bits seront les suivants

0 0 0 0 0 0 

Nous aurions 6 fils de travail sur le tableau d'entrée. Supposons que le tableau d'entrée final sera

1 9 3 9 2 5 

donc le tableau finale de bits serait:

0 0 0 0 1 0 

Je ne veux pas utiliser un tableau de bool parce que chacune des valeurs prendra 1 octet de mémoire ce qui est beaucoup car je ne veux travailler qu'avec des bits.

Est-il possible de réaliser quelque chose comme ça?

J'ai pensé à créer un tableau char où chaque valeur du tableau aura 8 bits. Cependant, que se passe-t-il si deux threads souhaitent changer différents bits du premier caractère du tableau? Ils devraient faire l'opération atomiquement même si le changement à l'intérieur du bit sera à des endroits différents. Donc, l'utilisation d'opérations atomiques perturbera probablement le parallélisme, et dans ce cas l'utilisation d'opérations atomiques n'est pas nécessaire, cela n'a aucun sens, mais devrait être utilisé en raison des contraintes d'utilisation d'un tableau de caractères plutôt que de quelque chose de plus spécialisé. comme un std::bitset

Merci d'avance.

Answer 1

Je fournis une réponse tardive à cette question pour l'enlever de la liste sans réponse.

Pour faire ce que vous voulez réaliser, vous pouvez définir un tableau de unsigned int s de longueur N/32, où N est la longueur des tableaux que vous comparez. Ensuite, vous pouvez utiliser atomicAdd pour écrire chaque bit d'un tel tableau, selon que deux éléments des tableaux sont égaux ou non.

Ci-dessous je fournir un exemple simple:

#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(); 
}