Come raccogliere in modo efficiente i dati dai thread in CUDA?

Question

Ho un'applicazione che risolve un sistema di equazioni in CUDA, so per certo che ogni thread può trovare fino a 4 soluzioni, ma come posso copiare nuovamente all'host?

Sto passando un enorme array con spazio sufficiente per tutti i thread archiviare 4 soluzioni (4 doppie per ciascuna soluzione) e un'altra con il numero di soluzioni per thread, tuttavia si tratta di una soluzione ingenua, ed è l'attuale collo di bottiglia del mio kernel.

Mi piace davvero ottimizzare questo. Il problema principale è concatenare un numero variabile di soluzioni per thread in un singolo array.

Answer 1

La funzionalità che stai cercando si chiama compattazione del flusso.

Probabilmente è necessario fornire un array che contenga spazio per 4 soluzioni per thread perché è probabile che il tentativo di memorizzare direttamente i risultati in una forma compatta crei così tante dipendenze tra i thread che le prestazioni hanno ottenuto nella possibilità di copiare meno dati ritornano all'host vengono persi da un tempo di esecuzione del kernel più lungo. L'eccezione a questo è se quasi tutti i thread non trovano soluzioni. In tal caso, potresti essere in grado di utilizzare un'operazione atomica per mantenere un indice in un array. Quindi, per ciascuna soluzione trovata, la si memorizzerebbe in una matrice in un indice e quindi si utilizzerà un'operazione atomica per aumentare l'indice. Penso che sarebbe sicuro usare atomicAdd() per questo. Prima di memorizzare un risultato, il thread usava atomicAdd() per aumentare l'indice di uno. atomicAdd() restituisce il vecchio valore e il thread può memorizzare il risultato usando il vecchio valore come indice.

Tuttavia, data una situazione più comune, dove c'è un buon numero di risultati, la soluzione migliore sarà eseguire un'operazione di compattazione come passaggio separato. Un modo per farlo è con thrust::copy_if. Vedi this question per qualche altro background.