Nelle prime tre parti di questa guida abbiamo visto come realizzare un programma CUDA funzionante, introducendo e sfruttando gli aspetti basilari del modello di programmazione: l'approccio alla GPU come un coprocessore dotato di memoria privata verso la quale spostare i dati e leggere i risultati, il concetto di kernel partizionato in una griglia di blocchi, ciascuno contenente i thread molto "leggeri" che eseguono effettivamente i calcoli, le parti fondamentali dell'architettura in termini di organizzazione in multiprocessori e la presenza di una shared memory in ciascuno di esso per migliorare la comunicazione tra i thread di un blocco e le prestazioni. Da questa parte in avanti, invece, inizieremo ad affrontare quegli aspetti che, una volta parallelizzato nella maniera corretta il nostro algoritmo, o parte di esso, per essere eseguito su architettura CUDA, ci permetteranno di raggiungere le massime performance.
Iniziamo subito parlando delle ottimizzazioni per gli accessi in memoria, e in particolare della coalescenza degli accessi stessi. Tale aspetto, in particolare con le architetture della generazione GeForce 8000 e 9000 (oggi disponibili anche come GeForce G/GT/GTS della serie 100 e sulle G210, GT 220, GTS 240 e GTS 250), è forse tra le più vincolanti ottimizzazioni da attuare per ottenere del codice efficiente.
Coalescenza degli accessi: cosa significa, e come ottenerla?
Rendere possibili degli accessi coalescenti in memoria significa riuscire ad accorpare più accessi (letture o scritture) in un'unica transazione del controller di memoria. Gli accessi in memoria dei chip grafici compatibili con CUDA devono seguire delle regole ben precise per minimizzare il numero di transazioni. L'unità a cui fare riferimento è quella dell'half-warp, cioè di 16 thread facenti parti dello stesso gruppo di 32 thread (warp) tenuti in esecuzione per ciascun ciclo di clock su un multiprocessore. A seconda della compute capability supportata dal dispositivo, cambiano anche le linee guida da seguire nell'organizzare gli accessi di un half warp per ottenere la loro coalescenza.
Per i chip grafici con compute capability 1.0 o 1.1, e cioè le soluzioni derivate da G80 e G92 sopra elencate, gli accessi di un half-warp sono coalescenti se leggono un'area contigua di memoria di:
- 64 byte – ogni thread legge una word: int, float, ...
- 128 byte – ogni thread legge una double-word: int2, float2, ...
- 256 byte – ogni thread legge una quad-word: int4, float4, ...
In più, devono essere rispettate le seguenti restrizioni:
- L'indirizzo iniziale di una regione deve essere multiplo della grandezza della regione
- Il k-esimo thread di un half-warp deve accedere al k-esimo elemento di un blocco (letto o scritto), gli accessi devono cioè essere perfettamente allineati tra i thread
C'è un'eccezione:
- Gli accessi rimangono coalescenti se alcuni thread non partecipano (cioè non eseguono la lettura o la scrittura in esame)
Riportiamo qua di seguito alcuni casi di accessi coalescenti o non coalescenti mostrati nel manuale di programmazione di CUDA. Negli esempi vengono mostrati solo accessi a dati da 32 bit per ciascun thread (64 byte per l'intero half warp). Lo stesso ragionamento vale per accessi a 64 bit e a 128 bit.
In questa immagine vediamo due accessi che si "incrociano" (a sinistra) e accessi che partono dall'indirizzo 132 invece che 128 (a destra). In entrambi i casi non sono rispettate le condizioni per la coalescenza degli accessi dal punto di vista dell'allineamento: nel primo caso, non c'è una corrispondenza "sequenziale" tra gli accessi, nel secondo non si parte da un multiplo della granularità della transazione (64 byte). L'esecuzione di accessi non coalescenti, con architetture che hanno compute capability 1.0 o 1.1, comporta l'esecuzione di ben 16 transazioni di memoria invece che 1.
La seconda immagine presenta invece due casi di accessi coalescenti. A sinistra, vediamo 16 accessi a 32 bit, che partono dall'indirizzo 128, con corrispondenza uno a uno. A destra, la situazione è simile, ma alcuni thread non eseguono accessi.
L'ultima immagine, presenta invece a sinistra un caso in cui solo parte degli accessi sono allineati, e ad un certo punto si disallineano, risultando quindi in accessi non coalescenti, e a destra il caso dove ciascun thread accede a 96 bit di dati (per un totale di 192 byte per tutto l'half warp). Non corrispondendo a nessuno dei casi di 32 bit (64 byte totali), 64 bit (128 byte totali) e 128 bit (256 byte totali) ci troviamo ancora in un caso di accessi non coalescenti, con 16 transazioni di memoria lanciate per ciascun accesso.
Nel caso di compute capability 1.2 e 1.3, quest'ultima supportata da tutti i chip derivati da GT200, le condizioni per ottenere la coalescenza degli accessi sono fortunatamente un po' più lasche. In particolare, una singola transazione di memoria è eseguita per un half warp se gli accessi di tutti i thread sono compresi all'interno dello stesso segmento di grandezza uguale a:
- 32 byte, se tutti i thread accedono a word di 8 bit
- 64 byte, se tutti i thread accedono a word di 16 bit
- 128 byte, se tutti i thread accedono word di 32 o 64 bit
La coalescenza degli accessi in una singola transazione, se essi risiedono nello stesso segmento come specificato, è ottenuta per tutti gli schemi di indirizzi richiesti dall'half warp, inclusi anche schemi dove più thread accedono allo stesso indirizzo. Se, invece, un half-warp indirizza parole in segmenti differenti, saranno eseguite tante transazioni quanti sono i segmenti indirizzati.
L'algoritmo che si occupa di lanciare le transazioni, cerca anche di ridurre, se possibile, la loro grandezza. Infatti, se in una transazione da 128 byte sono contenuti solo nella metà alta o bassa del segmento (per esempio, dati tutti a 32 bit allineati, ma anche pattern più complessi ma in un blocco di dimensioni inferiori), allora la grandezza della transazione verrà ridotta a 64 byte. Lo stesso accade nel caso di transazioni a 64 byte che possono essere contenute in transazioni da 32 byte, anche quando esse siano già il risultato di una riduzione da una transazione a 128 byte. La suddivisione in segmenti parte, come immaginabile, da multipli interi della grandezza del segmento stesso. Dunque, un segmento a 128 byte parte da 0, 128, 256, ... e, nel caso della riduzione della gandezza della transazione a 64 byte, i dati devo rientrare nella sua metà alta o bassa, dunque di fatto partire da multipli interi di 64.
L'immagine riportata mostra la situazione nel caso degli accessi in memoria globale con schede dotate di chip della famiglia di GT200. Gli accessi sono a tutti su dati di 32 bit ciascuno. Nel primo caso, vediamo una serie di accessi, non necessariamente 1 a 1, che però rientrano all'interno dello stesso blocco a 64 byte. Nel secondo caso, vediamo degli accessi, non allineati rispetto all'inizio del blocco, che però risiedono tutti all'interno dello stesso blocco da 128 byte. Nel terzo caso, infine, la transazione è allineata partendo prima dell'indirizzo 128, e data la grandezza dati di 32 bit, sarebbero richieste 2 transazioni da 128 byte. Per la parte prima dell'indirizzo 128, però, i dati si trovano nei 64 byte bassi del segmento, e nei rispettivi 32 byte bassi di quest'ultimo (partono dopo l'indirizzo 96). Dunque, viene eseguita solo una transazione a 32 byte. Per la parte di dati collocati dopo l'indirizzo 128, abbiamo i dati tutti contenuti prima dell'indirizzo 192, dunque compresi nei 64 byte alti. Di conseguenza, viene lanciata una transazione a 64 byte.
Coalescenza: l'esempio del float3
Uno dei modi per rendere coalescenti accessi che apparentemente non lo sono è quello di usare la shared memory come memoria di lavoro. Infatti, per la shared memory non ci sono problematiche legate alla coalescenza degli accessi, ed è per esempio possibile fare in modo di caricare i dati da memoria globale in maniera da minimizzare le transazioni, eseguire quelle operazioni che avrebbero reso gli accessi non coalescenti dalla shared memory, e poi eseguire le scritture leggendo da shared memory, ancora massimizzando la coalescenza degli accessi.
Accenniamo al fatto che, comunque, la shared memory, pur non necessitando di coalescenza degli accessi, può presentare problematiche di conflitto tra banchi di memoria. Esse però influenzano le prestazioni in maniera sicuramente inferiore rispetto al numero di transazioni di memoria eseguite e verrano ad ogni modo affrontate in una parte successiva del tutorial.
Passiamo ora ad un esempio pratico che ci mostrerà come trarre effettivo vantaggio della shared memory per rendere gli accessi in memoria globale coalescenti in un caso limite per le architetture con compute capability 1.0 e 1.1. Prendiamo il caso del float3, che come evidenziato precedentemente, con dati di 96 bit (32 bit per elemento del vettore), non permette in alcun modo di ottenere accessi coalescenti se ciascun thread legge un intero vettore.
Il codice base del kernel, che legge un vettore, incrementa in maniera differente i vari campi e infine lo riscrive in memoria, è il seguente:
__global__ void accessFloat3(float3 *d_in, float3 d_out)
{
int index = blockIdx.x * blockDim.x + threadIdx.x;
float3 a = d_in[index];
a.x += 2;
a.y += 3;
a.z += 4;
d_out[index] = a;
}
La figura mostra come vengono eseguiti gli accessi (alla prima lettura), che si distribuiscono, per lo stesso half warp, su tre blocchi contigui da 64 byte. Questa operazione, con le schede di compute capability 1.0 / 1.1 costa 16 transazioni.
L'idea è quella di fare in modo che ciascun thread legga sempre 3 campi, ma stavolta di vettori differenti, con uno spiazzamento dato dal numero stesso di thread.
La figura mostra come gli accessi vengono fatti adesso su una sequenza 1 a 1 di valori a 32 bit in tre cicli di lettura partendo da 0, 256 e 512.
Il codice diventa il seguente:
_global__ void accessInt3Shared(float *g_in, float *g_out)
{
int index = 3 * blockIdx.x * blockDim.x + threadIdx.x;
__shared__ float s_data[256*3];
s_data[threadIdx.x] = g_in[index];
s_data[threadIdx.x+256] = g_in[index+256];
s_data[threadIdx.x+512] = g_in[index+512];
__syncthreads();
float3 a = ((float3*)s_data)[threadIdx.x];
a.x += 2;
a.y += 3;
a.z += 4;
((float3*)s_data)[threadIdx.x] = a;
__syncthreads();
g_out[index] = s_data[threadIdx.x];
g_out[index+256] = s_data[threadIdx.x+256];
g_out[index+512] = s_data[threadIdx.x+512];
}
Dove notiamo :
- L'allocazione della shared memory per 256 * 3 elementi float (32 bit).
- La lettura da parte di ciascun thread dell'elemento all'indirizzo indicizzato dal proprio indice di thread più lo spiazzamento dato dal blocco corrente, moltiplicato per 3, dello stesso ad un offset di 256 elementi e ad un offset di 512 elementi, con le rispettive scritture nella shared memory.
- La synchthread, che si assicura che le operazioni di scrittura in memoria shared siano terminate prima di procedere con la computazione. Questo diventa necessario perchè i thread stanno leggendo e scrivendo dati diversi da quelli sui quali poi opereranno.
- L'operazione di lettura dei dati da parte di un thread in shared memory, eseguita castando i valori all'indice corrente a float3. Vengono dunque letti dallo stesso thread 3 valori float consecutivi.
- La parte di elaborazione, invariata.
- La scrittura dei dati elaborati, eseguita ancora tramite un casting.
- La syncthread che si assicura la scrittura dei dati in shared memory prima di passare alla scrittura dei risultati in memoria globale.
- La scrittura dei dati in memoria globale, esattamente inversa all'operazione di lettura.
Gli effetti di questa ottimizzazione sono molto importanti a livello di prestazioni, in quanto su array di 1 milione di float 3 possono portare ad almeno un ordine di grandezza di miglioramento (da circa 3300 a 330 millisecondi).
In generale, quando si ha a che fare con delle strutture, l'idea è quella di cercare di privilegiare un'organizzazione dei dati in memoria globale come struttura di array (SoA - dove, cioè, nello spazio di memoria linearizzato tutti i campi dello stesso tipo delle strutture sono uno di seguito all'altro), piuttosto che come array di strutture (AoS - dove, invece, sono le strutture ad assere sequenziali). Se le SoA non sono possibili, allora o si utilizza la shared memory come mostrato per il caso del float 3, o si può provare a forzare l'allineamento con la primitiva __align:
__align(X) dove X è 4, 8 o 16.
