Perché preoccuparsi di sapere su CUDA Warps?

Question

Ho GeForce GTX460 SE, quindi è: 6 SM x 48 core CUDA = 288 core CUDA. È noto che in un Warp contiene 32 thread e che in un blocco contemporaneamente (alla volta) può essere eseguito solo un Warp. Cioè, in un singolo multiprocessore (SM) è possibile eseguire simultaneamente un solo blocco, un Warp e solo 32 thread, anche se sono disponibili 48 core?

E inoltre, un esempio per distribuire Thread e Block in calcestruzzo può essere utilizzato threadIdx.x e blockIdx.x. Per assegnarli usa il kernel < < Blocchi, Thread >>>(). Ma come allocare un numero specifico di Warp-s e distribuirli, e se non è possibile allora perché preoccuparsi di sapere su Warps?

Answer 1

la situazione è un po 'più complicato di ciò che si descrive.

Le unità ALU (core), carico/archivio (LD/ST) e unità di funzione speciale (SFU) (verde nell'immagine) sono unità pipeline. Mantengono i risultati di molti calcoli o operazioni contemporaneamente, in varie fasi di completamento. Quindi, in un ciclo, possono accettare una nuova operazione e fornire i risultati di un'altra operazione avviata molto tempo fa (circa 20 cicli per le ALU, se non ricordo male). Quindi, un singolo SM in teoria ha risorse per elaborare contemporaneamente 48 * 20 cicli = 960 operazioni ALU, che è 960/32 thread per warp = 30 warps. Inoltre, può elaborare le operazioni LD/ST e le operazioni SFU a prescindere dalla loro latenza e velocità effettiva.

Gli schedulatori di curvatura (giallo nell'immagine) possono pianificare 2 * 32 thread per warp = 64 thread per le condotte per ciclo. Quindi questo è il numero di risultati che possono essere ottenuti per orologio. Quindi, dato che ci sono un mix di risorse di calcolo, 48 core, 16 LD/ST, 8 SFU, ognuno con latenze differenti, un mix di orditi viene elaborato nello stesso momento. Ad ogni ciclo, gli schedulatori di warp cercano di "accoppiare" due orditi per pianificare, per massimizzare l'utilizzo dell'SM.

Gli schedulatori di ordito possono emettere deformazioni sia da blocchi diversi, sia da luoghi diversi nello stesso blocco, se le istruzioni sono indipendenti. Quindi, gli orditi da più blocchi possono essere elaborati allo stesso tempo.

In aggiunta alla complessità, le deformazioni che eseguono istruzioni per le quali sono presenti meno di 32 risorse, devono essere inviate più volte per tutti i thread da servire. Ad esempio, ci sono 8 SFU, quindi significa che una distorsione contenente un'istruzione che richiede le SFU deve essere programmata 4 volte.

Questa descrizione è semplificata. Esistono anche altre restrizioni che determinano il modo in cui la GPU pianifica il lavoro. Puoi trovare ulteriori informazioni cercando nel web "architettura fermi".

Quindi, venendo alla tua domanda attuale,

perché preoccuparsi di conoscere Fili di ordito?

Conoscere il numero di thread in un ordito e prenderlo in considerazione diventa importante quando si tenta di ottimizzare le prestazioni del proprio algoritmo.Se non si seguono queste regole, si perde prestazioni:

• Nell'invocazione kernel, <<<Blocks, Threads>>>, tenta di scegliere un numero di thread che divide in modo uniforme con il numero di thread in una distorsione. Se non lo fai, si finisce con l'avvio di un blocco che contiene thread inattivi.

• nel kernel, cercare di avere ogni thread in una distorsione seguono lo stesso percorso di codice. Se non lo fai, ottieni quella che si chiama divergenza di ordito. Ciò accade perché la GPU deve eseguire l'intera curvatura attraverso ciascuno dei percorsi di codice divergenti.

• nel kernel, cercare di avere ogni thread in un carico di ordito e memorizzare i dati in modelli specifici. Ad esempio, fare in modo che i thread in un warp accedano alle parole consecutive a 32 bit nella memoria globale.

Answer 2

sono fili raggruppati in fili di ordito necessariamente nell'ordine, 1 - 32, 33 - 64 ...?

Sì, il modello di programmazione garantisce che i thread siano raggruppati in orditi in quello specifico ordine.

Come semplice esempio di ottimizzazione dei percorsi di codice divergenti può essere utilizzata la separazione di tutti i thread nel blocco in gruppi di 32 thread? Ad esempio: switch (threadIdx.s/32) {caso 0:/* 1 warp */break; caso 1:/* 2 ordito */rottura;/* Etc * /}

Esattamente :)

Quanti byte devono essere letti in una sola volta per singolo ordito: 4 byte * 32 Fili, 8 byte * 32 Filati o 16 byte * 32 Le discussioni? Per quanto ne so, l'unica transazione per la memoria globale in una volta riceve 128 byte.

Sì, le transazioni nella memoria globale sono 128 byte. Quindi, se ogni thread legge una parola a 32 bit da indirizzi consecutivi (probabilmente devono essere allineati a 128 byte), tutti i thread nel warp possono essere assistiti con una singola transazione (4 byte * 32 thread = 128 byte). Se ogni thread legge più byte, o se gli indirizzi non sono consecutivi, è necessario emettere più transazioni (con transazioni separate per ciascuna linea a 128 byte separata che viene toccata).

questo è descritto nel Manuale di programmazione CUDA 4.2, la sezione F.4.2, "Global Memory". C'è anche un blurb in là che dice che la situazione è diversa con i dati che vengono memorizzati nella cache solo in L2, in quanto la cache L2 ha linee di cache a 32 byte. Non so come fare in modo che i dati vengano memorizzati nella cache solo in L2 o in quante transazioni si finiscono.