Streaming Multiprocessor, Blocchi e Fili (CUDA)

Question

Qual è il rapporto tra un nucleo CUDA, un multiprocessore in streaming e il modello CUDA di blocchi e le discussioni?

Quello che viene mappato a ciò che è e ciò che parallelized e come? e ciò che è più efficiente, massimizzare il numero di blocchi o il numero di fili?

---

La mia comprensione attuale è che ci sono 8 core CUDA per multiprocessore. e che ogni nucleo cuda sarà in grado di eseguire un blocco cuda alla volta. e tutte le filettature in quel blocco vengono eseguite serialmente in quel particolare nucleo.

È corretto?

Answer 1

Il layout filo / blocco è descritto in dettaglio nel CUDA Programming guide . In particolare, il capitolo 4 stati:

  

L'architettura CUDA è costruito attorno ad una matrice scalabile di multithreaded Streaming Multiprocessor (SM). Quando un programma CUDA della CPU host invoca una griglia kernel, i blocchi della griglia vengono enumerate e distribuiti a multiprocessori con capacità di esecuzione disponibili. I fili di un blocco di filettatura vengono eseguiti contemporaneamente su un multiprocessore, e più blocchi di filettatura possono eseguire simultaneamente su un multiprocessore. Come blocchi di filettatura terminano, nuovi blocchi vengono lanciati sui multiprocessori lasciati liberi.

Ogni SM contiene 8 CUDA core, ed in qualsiasi momento stanno eseguendo un singolo filo di ordito di 32 fili - così prende 4 cicli di clock per emettere una singola istruzione per tutta ordito. Si può supporre che le discussioni in un dato ordito eseguire in lock-step, ma per sincronizzare tutta orditi, è necessario utilizzare __syncthreads().

Answer 2

Per la GTX 970 ci sono 13 Streaming Multiprocessor (SM) con 128 Cuda core ciascuno. CUDA Cores sono anche chiamati Stream Processors (SP).

È possibile definire le griglie che mappa blocchi per la GPU.

È possibile definire blocchi che mappano le discussioni a Stream Processors (128 Cuda core per SM).

Un'ordito è sempre costituito da 32 fili e tutte fili di ordito vengono eseguiti simulaneously.

Per utilizzare la piena potenza possibile di una GPU è necessario molto di più thread per SM rispetto al SM ha SP. Per ogni capacità di elaborazione c'è un certo numero di fili che possono risiedere in uno SM alla volta. Tutti i blocchi definiti sono in coda e attendere una SM di avere le risorse (numero di SP gratuito), quindi viene caricato. La SM inizia ad eseguire orditi. Poiché un ordito ha solo 32 fili e SM ha per esempio 128 SP un SM può eseguire 4 fili di ordito in un dato momento. La cosa è se i fili fanno accesso alla memoria del filo bloccherà fino alla sua richiesta di memoria è soddisfatta. In numeri: Un calcolo aritmetico sulla SP ha una latenza di 18-22 cicli mentre un accesso di memoria non-cache globale può richiedere fino a 300-400 cicli. Ciò significa che se i fili di ordito uno attendono dati solo un sottoinsieme delle 128 SP funzionerebbero. Perciò lo scheduler interruttori per eseguire un'altra ordito se disponibile. E se questo blocca ordito esegue il successivo e così via. Questo concetto è chiamato nasconde latenza. Il numero dei fili di ordito e la dimensione del blocco determinano la capienza (da quanti orditi la SM può scegliere di eseguire). Se l'occupazione è alto è più difficile che non c'è lavoro per le SP.

La tua affermazione che ciascun core CUDA eseguirà un blocco alla volta è sbagliata. Se si parla di Streaming Multiprocessor possono eseguire orditi da ogni thread che risiedono nella SM. Se un blocco ha una dimensione di 256 fili e tuoi allowes GPU 2048 discussioni a residente a SM ogni SM avrebbe 8 isolati residenti da cui il SM può scegliere orditi da eseguire. Tutte le filettature dei orditi eseguiti sono eseguite in parallelo.

Ci trovi numeri per le diverse capacità di elaborazione e GPU Architetture qui: https://en.wikipedia.org/wiki/CUDA#Limitations

È possibile scaricare un foglio di calcolo di occupazione da Nvidia residenza foglio di calcolo (da Nvidia) .

Answer 3

Il Distributore Compute lavoro sarà pianificare un blocco di filettatura (CTA) su un SM solo se la SM ha risorse sufficienti per il blocco filo (di memoria condivisa, orditi, registri, barriere, ...). risorse a livello di blocco Discussione tale memoria condivisa sono allocati. Il allocare crea orditi sufficienti per tutte le filettature del blocco di filettatura. Il gestore delle risorse alloca orditi utilizzando Round Robin alle sub-partizioni SM. Ogni sottopartizione SM contiene uno scheduler ordito, file di registro, e unità di esecuzione. Una volta che un ordito viene assegnato un sottopartizione rimarrà sul sottopartizione fino al completamento o viene anticipata dalla un cambio di contesto (architettura Pascal). Sul cambio di contesto ripristinare l'ordito verrà ripristinato lo stesso SM stesso ordito-id.

Quando tutte le discussioni in ordito hanno completato le attese ordito scheduler per tutte le istruzioni in circolazione emesse dalla ordito per completare e poi il manager delle risorse rilascia le risorse di livello ordito che comprendono warp-id e registrare file.

Quando tutti orditi in un blocco thread completo risorse livello di blocco poi vengono rilasciati e le notifichi SM il distributore Compute lavoro che il blocco è stato completato.

Una volta che un ordito è assegnato a un sottopartizione e tutte le risorse vengono allocate l'ordito è considerato senso attivo che lo scheduler ordito è attivamente monitoraggio dello stato dell'ordito. Ad ogni ciclo dell'ordito Scheduler determinare quali orditi attivi sono in stallo e che sono idonei per emettere un'istruzione. Lo scheduler ordito raccoglie la più alta priorità ammissibili di ordito e problemi 1-2 istruzioni consecutivi dalla curvatura. Le regole per dual-issue sono specifici per ciascuna architettura. Se un ordito emette un carico di memoria può continuare a istruzioni indipendenti eseguite fino a raggiungere un'istruzione dipendente. L'ordito sarà poi riferire in fase di stallo fino al completamento del carico. Lo stesso vale per le istruzioni matematiche dipendenti. L'architettura SM è progettata per nascondere sia ALU e la latenza di memoria passando per ciclo tra orditi.

Questa risposta non utilizza il nucleo termine CUDA come Questo introduce un modello mentale errato. core CUDA sono pipeline singola precisione floating point unità di esecuzione / interi. La latenza tasso di emissione e di dipendenza è specifico per ogni architettura. Ogni sottopartizione SM e SM ha altre unità di esecuzione comprende unità di carico / magazzini, in virgola mobile a doppia precisione, in virgola mobile mezza precisione, unità ramo, ecc.

Per ottimizzare le prestazioni sviluppatore deve capire il trade off di blocchi vs orditi vs registri / thread.

L'occupazione termine è il rapporto tra orditi attivi per orditi massimi di un SM. Kepler - architettura Pascal (tranne GP100) hanno 4 schedulatori ordito per SM. Il numero minimo di orditi per SM dovrebbe essere almeno uguale al numero di schedulatori ordito. Se l'architettura ha un'esecuzione latenza dipendente di 6 cicli (Maxwell e Pascal) allora si avrebbe bisogno almeno 6 orditi per scheduler che è 24 per SM (occupazione 24/64 = 37,5%) per coprire la latenza. Se i fili hanno un livello di istruzione parallelismo allora questo potrebbe essere ridotta. Quasi tutti i noccioli emettere istruzioni latenza variabili come carichi di memoria che possono prendere 80-1000 cicli. Ciò richiede orditi più attivi per ordito scheduler per nascondere la latenza. Per ogni kernel c'è un trade off tra il numero di punti orditi e altre risorse come la memoria condivisa o registri in modo ottimizzazione in caso di occupazione del 100% non è consigliato come qualche altro sacrificio sarà probabilmente fatta. Il profiler CUDA può aiutare a identificare tasso di istruzioni problema, di occupazione, e le ragioni di stallo al fine di aiutare lo sviluppatore a determinare l'equilibrio.

La dimensione di un blocco di filettatura può influire sulle prestazioni. Se il kernel ha grandi blocchi e usa le barriere di sincronizzazione poi bancarelle barriera può essere un venire motivi di stallo. Questo può essere alleviato riducendo laorditi per blocco filo.

Answer 4

Ci sono molteplici in streaming multiprocessore su un solo dispositivo.
Un SM può contenere più blocchi. Ogni blocco può contenere più thread.
Un SM hanno più core CUDA (come sviluppatore, non si dovrebbe preoccuparsi di questo perché si astrae dalla ordito), che lavoreranno su thread. SM lavorare sempre in ordito di fili (sempre 32). Una curvatura sarà solo lavoro sul filo dello stesso blocco.
SM e bloccare entrambi hanno limiti sul numero di filo, il numero di registro e memoria condivisa.