Perché la CPU è più veloce della GPU per piccoli input?

Question

Ho riscontrato che la CPU viene eseguita più velocemente della GPU per piccole dimensioni di input. Perchè è questo? Preparazione, trasferimento dei dati o cosa?

Ad esempio per il kernel e la funzione CPU (codice CUDA):

__global__ void squareGPU(float* d_in, float* d_out, unsigned int N) {
    unsigned int lid = threadIdx.x;
    unsigned int gid = blockIdx.x*blockDim.x+lid;
    if(gid < N) {
        d_out[gid] = d_in[gid]*d_in[gid]; 
    }
}

void squareCPU(float* d_in, float* d_out, unsigned int N) {
    for(unsigned int i = 0; i < N; i++) {
        d_out[i] = d_in[i]*d_in[i]; 
    }
}

Eseguendo queste funzioni 100 volte su un array di 5000 float a 32 bit, ottengo quanto segue utilizzando un piccolo programma di test

Size of array:
5000
Block size:
256

You chose N=5000 and block size: 256

Total time for GPU: 403 microseconds (0.40ms)
Total time for CPU: 137 microseconds (0.14ms)

Aumentando la dimensione dell'array a 1000000, ottengo:

Size of array:
1000000
Block size:
256

You chose N=1000000 and block size: 256

Total time for GPU: 1777 microseconds (1.78ms)
Total time for CPU: 48339 microseconds (48.34ms)

Non sto includendo il tempo utilizzato per trasferire i dati tra host e dispositivo (e viceversa), infatti, ecco la parte rilevante della mia procedura di test:

gettimeofday(&t_start, NULL);

for(int i = 0; i < 100; i++) {
    squareGPU<<< num_blocks, block_size>>>(d_in, d_out, N);
} cudaDeviceSynchronize();

gettimeofday(&t_end, NULL);

Dopo aver scelto la dimensione del blocco, calcolo il numero di blocchi relativamente alla dimensione dell'array: unsigned int num_blocks = ((array_size + (block_size-1)) / block_size);

Answer 1

Rispondere alla domanda generale del confronto delle prestazioni tra CPU e GPU è abbastanza complicato e generalmente implica la considerazione di almeno 3 o 4 diversi fattori a cui posso pensare. Tuttavia, hai semplificato un po 'il problema isolando la tua misurazione dai calcoli effettivi, al contrario del trasferimento dei dati o dell' "operazione completa".

In questo caso, ci sono probabilmente almeno 2 cose da considerare:

1. Overhead di avvio del kernel - L'avvio di un kernel su una GPU comporta un overhead di costo fisso "approssimativo", solitamente compreso tra 5 e 50 microsecondi, per avvio del kernel. Ciò significa che se si ridimensiona la quantità di lavoro in modo che la CPU possa farlo in meno di quella quantità di tempo, non è possibile che la GPU possa essere più veloce. Anche al di sopra di quel livello, c'è una funzione lineare che descrive quel modello di overhead, che sono sicuro che puoi elaborare se lo desideri, per confrontare le prestazioni della CPU rispetto alla GPU in presenza di un overhead di costo fisso. Quando si confrontano piccoli casi di test, questo è un fattore importante da considerare, tuttavia la mia ipotesi è che poiché la maggior parte dei tempi del caso di test sono ben al di sopra di 50 microsecondi, possiamo tranquillamente "ignorare" questo fattore, come approssimazione.

2. Le prestazioni / capacità effettive della CPU effettiva rispetto alla GPU effettiva. Questo è generalmente difficile da modellare, dipende dall'hardware specifico che stai utilizzando e non hai fornito queste informazioni. Tuttavia possiamo comunque fare alcune osservazioni e alcune congetture, espandendole nella prossima sezione, sulla base dei dati che ci hai fornito.

I tuoi due casi riguardano una quantità totale di lavoro descritta da N, considerando N=5000e N=1000000. Costruire un piccolo grafico:

      N  |  CPU time    |  GPU time
   5000  |    137       |  403
1000000  |  48339       | 1777

Quindi vediamo che nel caso della CPU, quando il lavoro è aumentato di un fattore 200, il tempo di esecuzione è aumentato di un fattore ~ ​​352, mentre nel caso della GPU, il tempo di esecuzione è aumentato di un fattore ~ ​​4.5. Avremo bisogno di spiegare entrambe queste "non linearità" per avere un'ipotesi ragionevole su cosa sta succedendo.

1. Effetti della cache: poiché stai eseguendo i tuoi casi di test 100 volte, le cache potrebbero avere un effetto. Nel caso della CPU, questa è la mia unica ipotesi sul motivo per cui non vedi una relazione lineare. Immagino che alle dimensioni molto ridotte, ci si trovi in ​​una cache "interna" della CPU, con 40 KB di dati "in vista". Passando alla dimensione maggiore, hai 8 MB di dati in vista, e sebbene questo probabilmente si adatti alla cache "esterna" della tua CPU, è possibile che non lo sia, e anche se lo fa, la cache esterna potrebbe produrre prestazioni complessive più lente di la cache interna. Immagino che questo sia il motivo per cui la CPU sembra peggiorare man mano che i dati diventano più grandi. La tua CPU viene influenzata in modo non lineare in modo negativomodo, dal set di dati più ampio. Nel caso della GPU, la cache esterna è al massimo di 6 MB (a meno che tu non stia utilizzando una GPU Ampere), quindi il tuo set di dati più grande non si adatta completamente alla cache esterna.

2. Effetti della saturazione della macchina: sia la CPU che la GPU possono essere completamente "caricate" o parzialmente, a seconda del carico di lavoro. Nel caso della CPU, immagino che tu non stia utilizzando alcun multi-threading, quindi il tuo codice CPU è limitato a un singolo core. (E la tua CPU ha quasi certamente più core disponibili.) Il tuo codice a thread singolo "saturerà" approssimativamente, cioè manterrà quel singolo core "occupato". Tuttavia la GPU ha molti core e immagino che il tuo test case più piccolo (che funzionerà a 5000 thread) saturerà solo parzialmente la tua GPU. Quello che voglio dire è che alcune delle risorse di elaborazione dei thread della GPU saranno inattive nel caso più piccolo (a meno che non ti capiti di essere in esecuzione sulla più piccola delle GPU). 5000 thread sono sufficienti per mantenere occupati 2 SM GPU, quindi se la tua GPU ha più di 2 SM,alcune delle sue risorse sono inattive durante il test case più piccolo, mentre il tuo test case più grande di un milione di thread è sufficiente per saturare, ovvero mantenere occupate tutte le risorse di elaborazione dei thread, su qualsiasi GPU CUDA corrente. L'effetto di ciò è che, sebbene la CPU non tragga alcun vantaggio da un test case più ampio (dovresti considerare l'utilizzo del multi-threading), la tua GPU probabilmente ne trarrà vantaggio. Il test case più grande consente alla tua GPU di lavorare di piùnella stessa quantità di tempo che sta impiegando il caso di test più piccolo. Pertanto la GPU beneficia in modo non lineare in modo positivo, dal carico di lavoro maggiore.

La GPU è anche in grado di mitigare meglio gli effetti della mancanza nella cache esterna, quando viene assegnato un carico di lavoro abbastanza grande. Questo è chiamato effetto di occultamento della latenza della GPU in presenza di un carico di lavoro parallelo "grande" e la CPU non ha (o non ha tanto) un meccanismo corrispondente. Quindi, a seconda della CPU e della GPU esatte, questo potrebbe essere un fattore aggiuntivo. Non intendo fornire un tutorial completo sull'occultamento della latenza qui, ma il concetto si basa parzialmente sull'elemento 2 sopra, quindi puoi raccogliere l'idea generale / trarne vantaggio.