Riduzioni parziali efficaci con matrici di elementi, offset e lunghezze di sottoliste

Question

Per la mia applicazione devo gestire un sacco di oggetti (diciamo int s) che vengono successivamente divisi e ordinati in contenitori più piccoli. A tal fine, posso conservare gli elementi in un singolo schiera continua

arr = {0,1,2,3,4,5,6,7,8,9,10,11,12,13,14...} 

e le informazioni relative alle benne (sottoliste) è dato dalla offset al primo elemento nella rispettiva secchio e le lunghezze della sottolista.

Così, per esempio, dato

offsets = {0,3,8,..} 
sublist_lengths = {3,5,2,...} 

comporterebbe seguenti cavalli:

0 1 2 || 3 4 5 6 7 || 8 9 || ... 

Quello che sto cercando è un modo un po 'generale ed efficiente per eseguire algoritmi, come riduzioni, sui bucket solo utilizzando i kernel personalizzati o la libreria thrust. Sommando i secchi dovrebbe dare:

3 || 25 || 17 || ... 

Quello che è venuta in mente:

• opzione 1: kernel personalizzati richiedono un bel po 'di ritocchi, le copie in memoria condivisa, scelta corretta di dimensioni di blocchi e griglie e una propria implementazione degli algoritmi, come scansione, riduzione, ecc. Inoltre, ogni singola operazione richiederebbe un proprio kernel personalizzato. In generale è chiaro a me come fare questo, ma dopo aver utilizzato thrust per l'ultimo paio di giorni ho l'impressione che ci potrebbe essere un modo più intelligente

• opzione 2: generare una serie di chiavi dalla l'offset ({0,0,0,1,1,1,1,1,2,2,3,...} nell'esempio precedente) e utilizzare thrust::reduce_by_key. Non mi piace la generazione di elenchi in più, però.

• opzione 3: Usa thrust::transform_iterator insieme thrust::counting_iterator per generare il sopra data lista chiave al volo. Sfortunatamente, non riesco a trovare un'implementazione che non richieda incrementi di indici alla lista di offset sul dispositivo e sconfigge il parallelismo.

Quale sarebbe il modo più corretto per implementarlo?

Answer 1

All'interno di Thrust, non riesco a pensare ad una soluzione migliore rispetto all'opzione 2. Le prestazioni non saranno terribili, ma certamente non sono ottimali.

La struttura dei dati presenta somiglianze con il formato Compressed Sparse Row (CSR) per la memorizzazione di matrici sparse, quindi è possibile utilizzare tecniche sviluppate per il calcolo di sparse matrix-vector multiplies (SpMV) per tali matrici se si desidera ottenere prestazioni migliori. Notare che la matrice "offset" del formato CSR ha lunghezza (N + 1) per una matrice con N righe (cioè i bucket nel tuo caso) dove l'ultimo valore di offset è la lunghezza di arr. Lo CSR SpMV code in Cusp è un po 'contorto, ma serve da buon punto di partenza per il tuo kernel. È sufficiente rimuovere qualsiasi riferimento a Aj o x dal codice e passare offsets e arr negli argomenti Ap e Av rispettivamente.

Answer 2

Non hai menzionato quanto sono grandi i secchi. Se i bucket sono abbastanza grandi, forse puoi cavartela copiando gli offset e le sottofinestre dell'host, eseguendo il iterazione su di essi e facendo una chiamata Thrust separata per ciascun bucket. Fermi può avere contemporaneamente 16 kernel in volo, quindi su quell'architettura potresti essere in grado di gestire secchi più piccoli e ottenere comunque un buon utilizzo.