blocchi, discussioni, deformabilità Dimensione

Question

C'è stata molta discussione su come scegliere i blocchi # & blockSize, ma mi manca ancora qualcosa. Molti dei miei preoccupazioni rispondere a questa domanda: How CUDA Blocks/Warps/Threads map onto CUDA Cores? (Per semplificare la discussione, c'è abbastanza perThread & memoria perBlock limiti di memoria non sono un problema qui..)

kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal); 

1) Per mantenere la SM occupato come possibile, Devo impostare nThreads su un multiplo di warpSize. Vero?

2) Un SM può eseguire solo un kernel alla volta. Questo è tutto ciò che HWcores di SM sta eseguendo solo kernelA. (Non alcuni HWcores che eseguono kernelA, mentre altri eseguono kernelB.) Quindi se ho solo un thread da eseguire, sto "sprecando" gli altri HWcores. Vero?

3) Se i problemi del warp-scheduler funzionano in unità di warpSize (32 thread) e ogni SM ha 32 HW, l'SM verrà utilizzato completamente. Cosa succede quando l'SM ha 48 HWcores? Come posso mantenere tutti i 48 core utilizzati al massimo quando lo scheduler sta lavorando in blocchi di 32? (Se il paragrafo precedente è vero, non sarebbe meglio se lo schedulatore emettesse lavoro in unità di dimensione HWcore?)

4) Sembra che lo schedulatore di warp accoda 2 attività alla volta. In modo che quando il kernel attualmente in esecuzione si blocca o blocchi, il secondo kernel viene scambiato. (Non è chiaro, ma suppongo che la coda qui sia più di 2 kernel profondi.) È corretto?

5) Se il mio HW ha un limite superiore di 512 thread per blocco (nThreadsMax), ciò non significa che il kernel con 512 thread funzionerà più velocemente su un blocco. (Anche in questo caso, mem non è un problema.) Ci sono buone probabilità di ottenere prestazioni migliori se diffondo il kernel a 512 thread su molti blocchi, non solo uno. Il blocco viene eseguito su uno o più SM. Vero?

5a) Penso che più piccolo è, meglio è, ma importa quanto piccolo faccio nBlocks? La domanda è, come scegliere il valore di nBlocks che è decente? (Non necessariamente ottimale.) Esiste un approccio matematico alla scelta di nBlocks o semplicemente trial-n-err.

Answer 1

Lasciami provare a rispondere alle tue domande una per una.

1. Che è corretto.
2. Cosa intendi esattamente per "HWcores"? La prima parte della tua affermazione è corretta.

3. Secondo lo NVIDIA Fermi Compute Architecture Whitepaper: "SM pianifica i thread in gruppi di 32 thread paralleli denominati warps. Ogni SM include due schedulatori di ordito e due unità di invio delle istruzioni, consentendo di emettere ed eseguire due orditi contemporaneamente. due distorsioni e invia un'istruzione da ogni distorsione a un gruppo di sedici core, sedici unità di carico/archivio o quattro SFU. Poiché gli orditi vengono eseguiti in modo indipendente, lo scheduler di Fermi non ha bisogno di verificare le dipendenze all'interno del flusso di istruzioni ".

   Inoltre, gli stati NVIDIA Keppler Architecture Whitepaper: "Il programma di scansione ordito quad di Kepler seleziona quattro orditi e due istruzioni indipendenti per curvatura possono essere inviate ogni ciclo."

   I nuclei "in eccesso" vengono quindi utilizzati programmando più di un ordito alla volta.

4. Lo scheduler di ordito pianifica gli orditi dello stesso kernel, non di kernel diversi.

5. Non proprio vero: ogni blocco è bloccato in un singolo SM, poiché è lì che risiede la sua memoria condivisa.
6. Questo è un problema complesso e dipende da come viene implementato il kernel. Si consiglia di dare un'occhiata al webinar nVidia Better Performance at Lower Occupancy di Vasily Volkov che spiega alcuni dei problemi più importanti. In primo luogo, tuttavia, suggerirei di scegliere il numero di thread per migliorare l'occupazione, utilizzando lo CUDA Occupancy Calculator.

Answer 2

1) Sì.

2) I dispositivi CC 2.0 - 3.0 possono eseguire contemporaneamente fino a 16 griglie. Ogni SM è limitato a 8 blocchi, quindi per raggiungere la piena concorrenza il dispositivo deve avere almeno 2 SM.

3) Sì, gli schedulatori di ordito selezionano ed emettono orditi al momento. Dimentica il concetto di core CUDA che sono irrilevanti. Per nascondere la latenza è necessario avere un parallelismo di istruzione elevato o un'occupazione elevata. Si consiglia di avere> 25% per CC 1.x e> 50% per CC> = 2.0. In generale CC 3.0 richiede un'occupazione superiore a 2.0 dispositivi a causa del raddoppio degli scheduler, ma solo un aumento del 33% di orditi per SM. L'esperimento Nsight VSE Issue Efficiency è il modo migliore per determinare se si dispone di warping sufficienti per nascondere le istruzioni e la latenza della memoria. Sfortunatamente, il Visual Profiler non ha questa metrica.

4) L'algoritmo del programma di curvatura non è documentato; tuttavia, non considera quale griglia ha originato il blocco di thread. Per i dispositivi CC 2.xe 3.0 il distributore di lavoro CUDA distribuirà tutti i blocchi da una griglia prima di distribuire i blocchi dalla griglia successiva; tuttavia, questo non è garantito dal modello di programmazione.

5) Per mantenere occupato il SM è necessario disporre di blocchi sufficienti per riempire il dispositivo. Dopodiché, assicurati di avere degli orditi sufficienti per raggiungere un'occupazione ragionevole. Ci sono sia vantaggi che svantaggi nell'usare blocchi di thread di grandi dimensioni. I blocchi di thread di grandi dimensioni in genere utilizzano meno cache di istruzioni e hanno impronte minori sulla cache; tuttavia, i blocchi di thread di grandi dimensioni si bloccano ai syncthreads (SM può diventare meno efficiente in quanto vi sono meno distorsioni tra cui scegliere) e tendono a mantenere istruzioni in esecuzione su unità di esecuzione simili. Raccomando di provare 128 o 256 thread per blocco thread per iniziare. Ci sono buone ragioni per blocchi di thread più grandi e più piccoli. 5a) Utilizzare il calcolatore di occupazione. Il prelievo di dimensioni troppo grandi di un blocco di filettatura spesso causa la limitazione dei registri. Il prelievo di dimensioni di blocco thread troppo piccole può risultare limitato dalla memoria condivisa o dagli 8 blocchi per limite SM.