Utilizzando variabile dispositivo più thread su CUDA

Question

sto giocando intorno con CUDA.

Al momento ho un problema. Sto testando una vasta gamma di particolari risposte, e quando ho la risposta, devo copiare i dati su un altro array.

Per esempio, la mia serie di prova di 5 elementi assomiglia a questo:
[] [] [V1] [] [] [v2]

Risultato deve apparire così:
[V1] [v2]

Il problema è come faccio a calcolare l'indirizzo del secondo array per memorizzare il risultato? Tutti gli elementi del primo array sono controllati in parallelo.

sto pensando di dichiarare una variabile int dispositivo di addr = 0. Ogni volta che trovo una risposta, ho incrementerà il addr . Ma io non sono sicuro di questo, perché significa che addr sono accessibili tramite più thread contemporaneamente. Sarà che causare problemi? O sarà il filo aspettare fino a quando un altro thread termina con quella variabile?

Answer 1

Non è così banale come sembra. Ho appena finito di implementare uno e vi posso dire che cosa avete bisogno leggere il scansione GPU Gems 3 Articolo in particolare il capitolo 39.3.1 Streaming compattazione .

Per implementare il proprio inizio dall'esempio LargeArrayScan nel SDK, che vi darà solo la scansione preliminare. Supponendo di avere la matrice di selezione nella memoria del dispositivo (un array di 1 e 0 significato 1- selezionare 0 - scartare), dev_selection_array dev_elements_array elementi da selezionare un dev_prescan_array e dev_result_array tutti dimensione N , quindi fare

prescan(dev_prescan_array,dev_selection_array, N);
scatter(dev_result_array, dev_prescan_array,
         dev_selection_array, dev_elements_array, N);

dove la dispersione è

 __global__ void scatter_kernel( T*dev_result_array, 
                   const T* dev_prescan_array, 
                   const T* dev_selection_array,
                   const T* dev_elements_array, std::size_t size){

unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= size) return;
if (dev_selection_array[idx] == 1){
    dev_result_array[dev_prescan_array[idx]] = dev_elements_array[idx];
}
}

per altri simpatica applicazione della scansione preliminare vede la carta Ble93

Buon divertimento!

Answer 2

Stai parlando di classico flusso di compattazione. In generale mi raccomando guardando spinta o CUDPP (quei link andare alla documentazione compattazione). Entrambi sono open source, se si vuole rotolare il proprio poi vorrei anche suggerire a guardare il campione 'scan' SDK.