Razionalizzare ciò che sta accadendo nel mio semplice kernel OpenCL per quanto riguarda la memoria globale

Question

const char programSource[] =
        "__kernel void vecAdd(__global int *a, __global int *b, __global int *c)"
        "{"
        "    int gid = get_global_id(0);"
        "for(int i=0; i<10; i++){"
        "    a[gid] = b[gid] + c[gid];}"
        "}";

Il kernel sopra è un'addizione vettoriale eseguita dieci volte per ciclo.Ho utilizzato la guida alla programmazione e l'overflow dello stack per capire come funziona la memoria globale, ma non riesco ancora a capire guardando il mio codice se sto accedendo alla memoria globale in modo corretto.Vi accedo in modo contiguo e indovino in modo allineato.La scheda carica blocchi da 128kb di memoria globale per gli array a, b e c?Quindi carica i blocchi da 128kb per ciascun array una volta per ogni 32 indici gid elaborati?(4*32=128) Sembra che non sto sprecando alcuna larghezza di banda della memoria globale, giusto?

A proposito, il profiler di calcolo mostra un'efficienza gld e gst di 1.00003, il che sembra strano, pensavo che sarebbe stato solo 1.0 se tutti i miei negozi e carichi fossero stati uniti.Come è sopra 1.0?

Answer 1

Sì, il tuo modello di accesso alla memoria è praticamente ottimale.Ogni semideformazione accede a 16 parole consecutive a 32 bit.Inoltre l'accesso è allineato a 64 byte, poiché i buffer stessi sono allineati e lo startindex per ogni halfwarp è un multiplo di 16.Quindi ogni halfwarp genererà una transazione da 64 byte.Quindi non dovresti sprecare larghezza di banda della memoria attraverso accessi non coalizzati.

Dato che hai chiesto esempi nella tua ultima domanda, modifichiamo questo codice per altri (modello di accesso meno ottimale (poiché il ciclo non fa davvero nulla, lo ignorerò):

kernel void vecAdd(global int* a, global int* b, global int* c)
{
   int gid = get_global_id(0);
   a[gid+1] = b[gid * 2] + c[gid * 32];
}

Inizialmente vediamo come funziona sull'hardware Compute 1.3 (GT200).

Per le scritture su a questo genererà un modello leggermente non ottimale (seguendo le semicurvature identificate dal loro intervallo id e dal modello di accesso corrispondente):

   gid  | addr. offset | accesses     | reasoning
  0- 15 |     4- 67    | 1x128B       | in aligned 128byte block
 16- 31 |    68-131    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
 32- 47 |   132-195    | 1x128B       | in aligned 128byte block
 48- 63 |   196-256    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access

Quindi, in pratica, stiamo sprecando circa la metà della nostra larghezza di banda (la larghezza di accesso inferiore a quella raddoppiata per gli dispari halfwarp non aiuta molto perché genera più accessi, il che non è più veloce dello spreco di più byte per così dire).

Per le letture da b i thread accedono solo agli elementi pari dell'array, quindi per ogni halfwarp tutti gli accessi si trovano in un blocco allineato di 128 byte (il primo elemento è al confine di 128B, poiché per quell'elemento il gid è un multiplo di 16=> l'indice è un multiplo di 32, per elementi da 4 byte, ciò significa che l'offset dell'indirizzo è un multiplo di 128B).Il modello di accesso si estende sull'intero blocco 128B, quindi questo eseguirà un trasferimento di 128B per ogni mezza curvatura, consumando nuovamente metà della larghezza di banda.

Le letture da c generano uno degli scenari peggiori, in cui ogni thread indicizza nel proprio blocco 128B, quindi ogni thread necessita del proprio trasferimento, che da un lato è un po' uno scenario di serializzazione (anche se non così grave come normalmente, poiché l'hardware dovrebbe essere in grado di sovrapporsi ai trasferimenti).Quel che è peggio è il fatto che questo trasferirà un blocco da 32B per ogni thread, sprecando 7/8 della larghezza di banda (accediamo a 4B/thread, 32B/4B=8, quindi viene utilizzato solo 1/8 della larghezza di banda).Poiché questo è il modello di accesso delle matrici ingenue, è altamente consigliabile farlo utilizzando la memoria locale (parlando per esperienza).

Calcolo 1.0 (G80)

Qui l'unico pattern che creerà un buon accesso è l'originale, tutti i pattern nell'esempio creeranno un accesso completamente non coalizzato, sprecando 7/8 della larghezza di banda (trasferimento/thread da 32B, vedere sopra).Per l'hardware G80 ogni accesso in cui l'ennesimo thread in un halfwarp non accede all'ennesimo elemento crea tali accessi non coalizzati

Calcolo 2.0 (Fermi)

Qui ogni accesso alla memoria crea 128B transazioni (tante quante necessarie per raccogliere tutti i dati, quindi 16x128B nel peggiore dei casi), tuttavia queste vengono memorizzate nella cache, rendendo meno ovvio dove verranno trasferiti i dati.Per il momento supponiamo che la cache sia abbastanza grande da contenere tutti i dati e non ci siano conflitti, quindi ogni linea di cache da 128B verrà trasferita al massimo una volta.Ipotizziamo inoltre un'esecuzione serializzata dei semicurvature, in modo da avere un'occupazione della cache deterministica.

Gli accessi a b trasferiranno comunque sempre 128B blocchi (nessun altro indice di thread nell'area di memoria corrispondente).L'accesso a c genererà 128B trasferimenti per thread (peggiore modello di accesso possibile).

Per gli accessi ad a è il seguente (trattandoli per il momento come reads):

   gid  | offset  | accesses | reasoning
  0- 15 |   4- 67 |  1x128B  | bringing 128B block to cache
 16- 31 |  68-131 |  1x128B  | offsets 68-127 already in cache, bring 128B for 128-131 to cache
 32- 47 | 132-195 |    -     | block already in cache from  last halfwarp
 48- 63 | 196-259 |  1x128B  | offsets 196-255 already in cache, bringing in 256-383

Quindi, per gli array di grandi dimensioni, gli accessi a un teoricamente non sprecheranno quasi alcuna larghezza di banda.Per questo esempio la realtà ovviamente non è altrettanto buona, poiché gli accessi a c distruggeranno abbastanza bene la cache

Per il profiler suppongo che le efficienze superiori a 1.0 siano semplicemente il risultato di inesattezze in virgola mobile.

Spero possa aiutare