Come ottimizzare il gioco della vita di Conway per CUDA?

Question

ho scritto questo kernel CUDA per la partita di Conway della vita:

__global__ void gameOfLife(float* returnBuffer, int width, int height) { 
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; 
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; 
    float p = tex2D(inputTex, x, y); 
    float neighbors = 0; 
    neighbors += tex2D(inputTex, x+1, y); 
    neighbors += tex2D(inputTex, x-1, y); 
    neighbors += tex2D(inputTex, x, y+1); 
    neighbors += tex2D(inputTex, x, y-1); 
    neighbors += tex2D(inputTex, x+1, y+1); 
    neighbors += tex2D(inputTex, x-1, y-1); 
    neighbors += tex2D(inputTex, x-1, y+1); 
    neighbors += tex2D(inputTex, x+1, y-1); 
    __syncthreads(); 
    float final = 0; 
    if(neighbors < 2) final = 0; 
    else if(neighbors > 3) final = 0; 
    else if(p != 0) final = 1; 
    else if(neighbors == 3) final = 1; 
    __syncthreads(); 
    returnBuffer[x + y*width] = final; 
} 

Cerco errori/ottimizzazioni. La programmazione parallela è abbastanza nuova per me e non sono sicuro di sapere come farlo nel modo giusto.

Il resto è una memcpy da un array di input all'input di texture 2DTex associato a un array CUDA. L'output è memcpyed dalla memoria globale all'host e quindi gestito.

Come si può vedere un thread si occupa di un singolo pixel. Non sono sicuro che sia il modo più veloce in quanto alcune fonti suggeriscono di fare una riga o più per thread. Se capisco correttamente, gli stessi NVidia dicono che più thread, meglio è. Mi piacerebbe consigli su questo da qualcuno con esperienza pratica.

Answer 1

I miei due centesimi.

Sembra che tutta la cosa sia limitata dalla latenza della comunicazione tra i multiprocessore e la memoria della GPU. Avete un codice che dovrebbe prendere qualcosa come 30-50 tick di clock da eseguire da solo, e genera almeno 3 accessi di memoria che richiedono più di 200 tick di clock se i dati richiesti non sono nella cache.

L'utilizzo della memoria texture è un buon modo per affrontarlo, ma non è necessariamente il modo ottimale.

Per lo meno, prova a fare 4 pixel alla volta (in orizzontale) per filo. È possibile accedere alla memoria globale a 128 byte alla volta (se si dispone di un curvatura che tenta di accedere a qualsiasi byte in un intervallo di 128 byte, si potrebbe anche inserire l'intera riga della cache quasi senza costi aggiuntivi). Dal momento che un warp è 32 thread, avere ogni thread funziona su 4 pixel dovrebbe essere efficiente.

Inoltre, si desidera che i pixel adiacenti verticalmente lavorati dallo stesso multiprocessore. Il motivo è che le file adiacenti condividono gli stessi dati di input. Se il pixel (x = 0, y = 0) è lavorato da un MP e il pixel (x = 0, y = 1) viene elaborato da un MP diverso, entrambi i MP devono emettere tre richieste di memoria globale ciascuna. Se entrambi sono lavorati dallo stesso MP e i risultati sono correttamente memorizzati nella cache (implicitamente o esplicitamente), è necessario solo un totale di quattro. Questo può essere fatto avendo ogni thread funzionante su diversi pixel verticali, o avendo blockDim.y> 1.

Più in generale, è probabile che si desideri che ogni warp a 32 thread carichi la quantità di memoria disponibile sul MP ​​(16-48 kb, o almeno un blocco 128x128), quindi elabora tutti i pixel all'interno quella finestra.

Su dispositivi di compatibilità di calcolo prima della 2.0, è consigliabile utilizzare la memoria condivisa. Sui dispositivi di compatibilità di calcolo 2.0 e 2.1, le funzionalità di caching sono molto migliorate, quindi la memoria globale potrebbe andare bene.

Alcuni risparmi non banali possono essere ottenuti assicurandosi che ogni warp acceda solo a due linee di cache in ogni riga orizzontale di pixel di input, anziché tre, come accadrebbe in un'implementazione ingenua che funziona su 4 pixel per thread, 32 thread per ordito.

Non c'è alcun motivo valido per utilizzare float come tipo di buffer. Non si finisce solo con quattro volte la larghezza di banda della memoria, ma il codice diventa inaffidabile e soggetto a errori. (Ad esempio, sei sicuro che if(neighbors == 3) funzioni correttamente, dal momento che stai confrontando un float e un intero?) Usa il char senza segno. Meglio ancora, usare uint8_t e typedef per significare unsigned char se non è definito.

Infine, non sottovalutare il valore della sperimentazione.Molto spesso le prestazioni del codice CUDA non possono essere facilmente spiegate dalla logica e devi ricorrere a parametri di regolazione e vedere cosa succede.

Answer 2

uno sguardo a questa discussione, abbiamo fatto assegnare di miglioramenti laggiù ...

http://forums.nvidia.com/index.php?showtopic=152757&st=60

Answer 3

TL; DR: vedi: http://golly.sourceforge.net

Il problema è che la maggior parte delle implementazioni CUDA seguono il idea cerebrale del conteggio manuale dei vicini. Questo è così lento e lento che qualsiasi implementazione di CPU seriale intelligente sarà più performante.

L'unico metodo ragionevole per eseguire calcoli GoL è l'utilizzo di una tabella di ricerca.
Le implementazioni attualmente più veloci su un utilizzo della CPU cercano un blocco quadrato 4x4 = 16 bit per vedere le future cellule 2x2 all'interno.

in questa configurazione le celle sono disposte in questo modo:

0xxxxxxxx //byte0 
1xxxxxxxx //byte1 
2 etc 
3 
4 
5 
6 
7 

Alcuni bit-shifting è impiegato per ottenere un blocco 4x4 di inserirsi in una parola e che la parola viene cercato utilizzando una tabella di ricerca. Le tabelle di ricerca contengono anche le parole, in questo modo è possibile memorizzare 4 diverse versioni del risultato nella tabella di ricerca, in modo da ridurre al minimo la quantità di bithifting necessaria per l'input e/o l'output.

Inoltre le diverse generazioni sono sfalsati, in modo da avere solo a guardare a 4 lastre vicini, invece di 9. Come così:

AAAAAAAA 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
      BBBBBBBB 
//odd generations (A) are 1 pixel above and to the right of B, 
//even generations (B) are 1 pixels below and to the left of A. 

Questo solo si traduce in una 1000x + speed-up rispetto ai sciocche implementazioni di conteggio.

Poi c'è l'ottimizzazione del calcolo non lastre che sono statici o hanno una periodicità di 2.

E poi c'è HashLife, ma questo è una bestia completamente diversa.
HashLife può generare modelli di vita in tempo O (log n), invece del tempo O (n) può essere eseguito dalle normali implementazioni. Ciò consente di calcolare la generazione: 6,366,548,773,467,669,985,195,496,000 (6 ottilioni) in pochi secondi.
Purtroppo Hashlife richiede la ricorsione, e quindi è difficile su CUDA.