Streaming multiprocessori, blocchi e thread (CUDA)

Question

Qual è la relazione tra un core CUDA, un multiprocessore di streaming e il modello CUDA di blocchi e thread?

Cosa viene mappato su cosa e cosa è parallelizzato e come? e cosa è più efficiente, massimizza il numero di blocchi o il numero di thread?

---

La mia attuale comprensione è che ci sono 8 cuda core per multiprocessore. e che ogni nucleo di cuda sarà in grado di eseguire un blocco di cuda alla volta. e tutti i thread in quel blocco sono eseguiti in serie in quel particolare core.

È corretto?

Answer 1

Il layout del thread/blocco è descritto in dettaglio nello CUDA programming guide. In particolare, il capitolo 4 riporta:

L'architettura CUDA è costruita attorno a una serie scalabile di multithreaded Streaming Multiprocessor (SM). Quando un programma CUDA sulla CPU host richiama una griglia del kernel, i blocchi della griglia vengono enumerati e distribuiti a multiprocessori con capacità di esecuzione disponibile. I thread di un blocco di thread vengono eseguiti contemporaneamente su un multiprocessore e più blocchi di thread possono essere eseguiti simultaneamente su un multiprocessore. Quando i blocchi di thread terminano, vengono avviati nuovi blocchi sui multiprocessori vuoti.

Ogni SM contiene 8 CUDA core, ed in qualsiasi momento che stanno eseguendo un singolo filo di ordito di 32 fili - quindi prende 4 cicli di clock per emettere una singola istruzione per tutta ordito. Puoi presumere che i thread in qualsiasi warp vengano eseguiti in lock-step, ma per sincronizzare tra gli orditi, devi usare __syncthreads().

Answer 2

Per la GTX 970 ci sono 13 Streaming Multiprocessor (SM) con 128 core Cuda ciascuno. I core Cuda sono anche chiamati stream processor (SP).

È possibile definire le griglie che mappano i blocchi alla GPU.

È possibile definire i blocchi che mappano i thread ai processori di flusso (i 128 Cuda Cores per SM).

Un ordito è sempre formato da 32 fili e tutti i fili di un ordito vengono eseguiti simultaneamente.

Per utilizzare tutta la potenza possibile di una GPU sono necessari molti più thread per SM rispetto a SM con SP. Per ogni capacità di calcolo esiste un certo numero di thread che possono risiedere in un SM alla volta. Tutti i blocchi che definisci sono in coda e attendi che un SM abbia le risorse (numero di SP liberi), quindi viene caricato. L'SM inizia a eseguire Warps. Poiché un Warp ha solo 32 Thread e un SM ha per esempio 128 SP un SM può eseguire 4 Warps in un dato momento. Il fatto è che se i thread accedono alla memoria, il thread si bloccherà fino a quando la sua richiesta di memoria non sarà soddisfatta. In numeri: un calcolo aritmetico sull'SP ha una latenza di 18-22 cicli mentre un accesso alla memoria globale non memorizzato nella cache può richiedere fino a 300-400 cicli. Ciò significa che se i thread di un ordito sono in attesa di dati, solo un sottoinsieme dei 128 SP funzionerebbe. Pertanto, il programmatore passa a eseguire un altro ordito, se disponibile. E se questo warp blocca, esegue il successivo e così via. Questo concetto è chiamato latency hiding. Il numero di orditi e la dimensione del blocco determinano l'occupazione (da quanti orditi l'SM può scegliere di eseguire). Se l'occupazione è elevata, è più improbabile che non ci sia lavoro per gli SP.

L'affermazione che ogni cuda core eseguirà un blocco alla volta è errata. Se parli di Streaming Multiprocessor, possono eseguire warp da tutti i thread che risiedono nell'SM.Se un blocco ha una dimensione di 256 thread e la GPU consente la permanenza di 2048 thread per SM, ogni SM avrà 8 blocchi residenti da cui SM può scegliere i warping da eseguire. Tutti i thread degli orditi eseguiti vengono eseguiti in parallelo.

a trovare i 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 di Nvidia Occupancy Calculation sheet (by 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 (memoria condivisa, fili di ordito, i registri, le barriere, ...). Le risorse a livello del blocco di thread vengono allocate come memoria condivisa. L'allocazione crea warping sufficienti per tutti i thread nel blocco thread. Il gestore risorse assegna il round robin warps alle sotto-partizioni SM. Ciascuna partizione SM contiene uno scheduler di warp, un file di registro e unità di esecuzione. Una volta assegnata una distorsione a una sottopartizione, rimarrà nella sottopartizione finché non viene completata o viene anticipata da un interruttore di contesto (architettura Pascal). Sul commutatore di contesto, il warp verrà ripristinato sullo stesso ID di warp stesso.

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 i Resource Manager rilascia le risorse a livello di ordito che comprendono ordito-id e registrare file.

Quando tutti gli orditi in un blocco di thread sono completi, le risorse a livello di blocco vengono rilasciate e SM comunica al Distributore di lavoro di calcolo che il blocco è stato completato.

Una volta assegnata una distorsione a una sottopartizione e assegnate tutte le risorse, l'ordito è considerato attivo, il che significa che lo scheduler di ordito sta attivamente monitorando lo stato del curvatura. Ad ogni ciclo, lo scheduler di ordito determina quali deformazioni attive sono bloccate e quali sono idonee a emettere un'istruzione. Lo scheduler del warp seleziona la warp idonea con la priorità più alta e rilascia 1-2 istruzioni consecutive dal warp. Le regole per la duplicazione sono specifiche per ogni architettura. Se un ordito emette un carico di memoria, può continuare a eseguire istruzioni indipendenti fino a quando non raggiunge un'istruzione dipendente. Il warp riporterà quindi in stallo fino al completamento del caricamento. Lo stesso vale per le istruzioni matematiche dipendenti. L'architettura SM è progettata per nascondere sia l'ALU che la latenza della memoria passando per ciclo tra gli orditi.

Questa risposta non utilizza il termine nucleo CUDA in quanto introduce un modello mentale errato. I core CUDA sono unità di precisione in virgola mobile/intero con pipeline a precisione singola. La frequenza di emissione e la latenza delle dipendenze sono specifiche per ogni architettura. Ciascuna SM subpartition e SM ha altre unità di esecuzione comprese unità load/store, unità a virgola mobile a precisione doppia, unità a virgola mobile di precisione, unità di derivazione, ecc.

Al fine di ottimizzare le prestazioni lo sviluppatore deve capire il trade off di blocchi contro deformazioni vs registri/thread.

Il termine occupazione è il rapporto tra gli orditi attivi e gli orditi massimi su una SM. Kepler - L'architettura Pascal (eccetto GP100) ha 4 programmatori di ordito per SM. Il numero minimo di orditi per SM dovrebbe essere almeno uguale al numero di schedulatori di ordito. Se l'architettura ha una latenza di esecuzione dipendente di 6 cicli (Maxwell e Pascal), allora occorrono almeno 6 warps per scheduler che sono 24 per SM (24/64 = 37,5% di occupazione) per coprire la latenza. Se i thread hanno parallelismo a livello di istruzioni, questo potrebbe essere ridotto. Quasi tutti i kernel emettono istruzioni di latenza variabile come carichi di memoria che possono richiedere 80-1000 cicli. Ciò richiede più warps attivi per lo scheduler di warp per nascondere la latenza. Per ogni kernel esiste un trade off point tra il numero di warp e altre risorse come la memoria condivisa oi registri, quindi l'ottimizzazione per il 100% di occupazione non è consigliata in quanto probabilmente verranno fatti altri sacrifici.Il profiler CUDA può aiutare a identificare il tasso di emissione delle istruzioni, l'occupazione e le ragioni di stallo al fine di aiutare lo sviluppatore a determinare tale equilibrio.

Le dimensioni di un blocco di filettatura possono influire sulle prestazioni. Se il kernel ha blocchi di grandi dimensioni e utilizza le barriere di sincronizzazione, le bancarelle di barriera possono essere un motivo di stallo. Questo può essere alleviato riducendo gli orditi per blocco di thread.