Ho caricato su Github .
È una semplice tabella hash per GPU in grado di gestire centinaia di milioni di inserimenti al secondo. Sul mio laptop con NVIDIA GTX 1060, il codice inserisce 64 milioni di coppie chiave-valore generate casualmente in circa 210 ms e rimuove 32 milioni di coppie in circa 64 ms.
Quindi, la velocità sul laptop è di circa 300 milioni di inserimenti/sec e 500 milioni di rimozioni/sec.
La tabella è scritta in CUDA, anche se la stessa metodologia può essere applicata a HLSL o GLSL. L'implementazione ha alcune limitazioni che garantiscono alte prestazioni sulla scheda grafica:
- Vengono elaborati solo chiavi e valori a 32 bit.
- La tabella hash ha una dimensione fissa.
- E questa dimensione deve essere una potenza di due.
È necessario riservare un marcatore delimitatore semplice per le chiavi e i valori (nel codice fornito è 0xffffffff).
Tabella hash senza blocchi
Nella tabella hash si utilizza l'indirizzamento aperto con , ovvero è semplicemente un array di coppie chiave-valore, che viene memorizzato in memoria e ha un'eccellente prestazione della cache. Questo non si può dire per il chaining, che implica la ricerca di un puntatore in una lista concatenata. La tabella hash è un semplice array che memorizza gli elementi KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
La dimensione della tabella è una potenza di due, non un numero primo, perché per applicare la maschera pow2/AND è suffetta una singola istruzione veloce, mentre l'operatore modulo è molto più lento. Questo è importante nel caso di sondatura lineare, poiché durante la ricerca lineare nella tabella, l'indice dello slot deve essere avvolto in ogni slot. E di conseguenza si aggiunge il costo dell'operazione modulo in ogni slot.
La tabella memorizza solo chiave e valore per ogni elemento, non l'hash della chiave. Poiché la tabella memorizza solo chiavi a 32 bit, l'hash viene calcolato molto rapidamente. Nel codice fornito viene utilizzato l'hash Murmur3, che esegue solo alcune rotazioni, XOR e moltiplicazioni.
Nella tabella hash viene utilizzata una tecnica di protezione contro i blocchi, che è indipendente dall'ordine di posizionamento in memoria. Anche se alcune operazioni di scrittura violano l'ordine di altre operazioni simili, la tabella hash manterrà comunque uno stato corretto. Ne parleremo più avanti. La tecnica funziona molto bene con le schede grafiche, in cui migliaia di thread vengono eseguiti in modo concorrente.
Le chiavi e i valori nella tabella hash vengono inizializzati a vuoti.
Il codice può essere modificato per gestire chiavi e valori a 64 bit. Le chiavi richiedono operazioni atomiche di lettura, scrittura e confronto con scambio (compare-and-swap). E per i valori sono necessarie operazioni atomiche di lettura e scrittura. Per fortuna, in CUDA le operazioni di lettura-scrittura per valori a 32 e 64 bit sono atomiche fintanto che sono allineate naturalmente (vedi ), e le schede grafiche moderne supportano operazioni atomiche di confronto con scambio a 64 bit. Ovviamente, passando ai 64 bit, le prestazioni potrebbero leggermente diminuire.
Stato della tabella hash
Ogni coppia chiave-valore nella tabella hash può avere uno dei quattro stati:
- La chiave e il valore sono vuoti. In questo stato, la tabella hash viene inizializzata.
- La chiave è stata scritta, ma il valore no. Se un altro thread legge i dati in quel momento, restituirà un valore vuoto. Questo è normale, accadrebbe lo stesso se un altro thread fosse stato elaborato un attimo prima, poiché stiamo parlando di una struttura dati concorrente.
- Sia la chiave che il valore sono stati scritti.
- Il valore è disponibile per altri thread, ma la chiave non lo è ancora. Questo può succedere perché il modello di programmazione in CUDA prevede un modello di memoria debolmente ordinato. Questo è normale, in qualsiasi evento la chiave è ancora vuota, anche se il valore non lo è più.
Un punto importante è che una volta che la chiave è stata scritta nello slot, non si sposta più - anche se la chiave viene rimossa, ne parleremo più avanti.
Il codice della tabella hash funziona anche con modelli di memoria debolmente ordinati, in cui non si conosce l'ordine delle letture e scritture in memoria. Quando esamineremo inserimenti, ricerche e cancellazioni nella tabella hash, ricordate che ogni coppia chiave-valore si trova in uno dei quattro stati descritti sopra.
Inserimento nella tabella hash
La funzione CUDA che inserisce coppie chiave-valore nella tabella hash appare così:
void gpu_hashtable_insert(KeyValue* hashtable, uint32_t key, uint32_t value)
{
uint32_t slot = hash(key);
while (true)
{
uint32_t prev = atomicCAS(&hashtable[slot].key, kEmpty, key);
if (prev == kEmpty || prev == key)
{
hashtable[slot].value = value;
break;
}
slot = (slot + 1) & (kHashTableCapacity-1);
}
}
Per inserire una chiave, il codice itera attraverso l'array della tabella hash partendo dall'hash della chiave da inserire. In ciascuno slot dell'array viene eseguita un'operazione atomica di confronto e scambio, in cui la chiave in quello slot viene confrontata con un valore vuoto. Se viene trovata una discrepanza, la chiave nello slot viene aggiornata con la chiave da inserire, e poi viene restituita la chiave originale dello slot. Se questa chiave originale era vuota o corrispondeva alla chiave da inserire, significa che il codice ha trovato uno slot adatto per l'inserimento e inserisce il valore nello slot.
Se in una singola chiamata del kernel gpu_hashtable_insert() ci sono più elementi con la stessa chiave, allora uno qualsiasi dei loro valori può essere scritto nello slot della chiave. Questo è considerato normale: una delle operazioni di scrittura della chiave-valore durante la chiamata avrà successo, ma poiché tutto ciò avviene in parallelo all'interno di più thread di esecuzione, non possiamo prevedere quale operazione di scrittura in memoria sarà l'ultima.
Ricerca nella tabella hash
Codice per la ricerca delle chiavi:
uint32_t gpu_hashtable_lookup(KeyValue* hashtable, uint32_t key)
{
uint32_t slot = hash(key);
while (true)
{
if (hashtable[slot].key == key)
{
return hashtable[slot].value;
}
if (hashtable[slot].key == kEmpty)
{
return kEmpty;
}
slot = (slot + 1) & (kHashTableCapacity - 1);
}
}
Per trovare il valore di una chiave memorizzata nella tabella, iteriamo attraverso l'array a partire dall'hash della chiave cercata. In ciascuno slot controlliamo se la chiave corrisponde a quella che stiamo cercando e, se sì, restituiamo il suo valore. Controlliamo anche se la chiave è vuota e, se è così, interrompiamo la ricerca.
Se non riusciamo a trovare la chiave, il codice restituisce un valore vuoto.
Tutte queste operazioni di ricerca possono essere eseguite in modo concorrente durante le operazioni di inserimento e cancellazione. Ogni coppia nella tabella avrà per il thread uno dei quattro stati sopra descritti.
Cancellazione nella tabella hash
Codice per la cancellazione delle chiavi:
void gpu_hashtable_delete(KeyValue* hashtable, uint32_t key, uint32_t value)
{
uint32_t slot = hash(key);
while (true)
{
if (hashtable[slot].key == key)
{
hashtable[slot].value = kEmpty;
return;
}
if (hashtable[slot].key == kEmpty)
{
return;
}
slot = (slot + 1) & (kHashTableCapacity - 1);
}
}
La rimozione di una chiave avviene in modo insolito: lasciamo la chiave nella tabella e contrassegniamo il suo valore (non la chiave stessa) come vuoto. Questo codice è molto simile a lookup(), tranne per il fatto che, quando viene trovata una corrispondenza per la chiave, ne rende il valore vuoto.
Come accennato in precedenza, una volta che la chiave è scritta nello slot, non si sposta più. Anche quando si rimuove un elemento dalla tabella, la chiave rimane al suo posto, il suo valore diventa semplicemente vuoto. Questo significa che non abbiamo bisogno di utilizzare un'operazione atomica di scrittura nello slot, perché non importa se il valore corrente è vuoto o meno: diventerà comunque vuoto.
Ridimensionamento della tabella hash
Il ridimensionamento della tabella hash può avvenire creando una tabella più grande e inserendo in essa gli elementi non vuoti dalla vecchia tabella. Non ho implementato questa funzionalità perché volevo mantenere il codice semplice. Inoltre, nelle applicazioni CUDA, l'allocazione della memoria viene spesso eseguita nel codice host anziché nel kernel CUDA.
Nell'articolo si discute di come modificare una tale struttura dati protetta da blocchi.
Concorrenza
Nei frammenti di codice sopra, le funzioni gpu_hashtable_insert(), _lookup() e _delete() gestiscono una coppia chiave-valore alla volta. Qui sotto, gpu_hashtable_insert(), _lookup() e _delete() gestiscono un array di coppie in parallelo, ogni coppia in un thread di esecuzione GPU separato:
// CPU code to invoke the CUDA kernel on the GPU
uint32_t threadblocksize = 1024;
uint32_t gridsize = (numkvs + threadblocksize - 1) / threadblocksize;
gpu_hashtable_insert_kernel<<<gridsize, threadblocksize>>>(hashtable, kvs, numkvs);
// GPU code to process numkvs key/values in parallel
void gpu_hashtable_insert_kernel(KeyValue* hashtable, const KeyValue* kvs, unsigned int numkvs)
{
unsigned int threadid = blockIdx.x*blockDim.x + threadIdx.x;
if (threadid < numkvs)
{
gpu_hashtable_insert(hashtable, kvs[threadid].key, kvs[threadid].value);
}
}
Una tabella hash con protezione da blocchi supporta inserti, ricerche e rimozioni concorrenti. Poiché le coppie chiave-valore si trovano sempre in uno dei quattro stati e le chiavi non si spostano, la tabella garantisce la correttezza anche con l'uso simultaneo di operazioni diverse.
Tuttavia, se stiamo elaborando parallelamente un batch di inserti e rimozioni, e se nell'array di ingressi ci sono chiavi duplicate, non possiamo prevedere quali coppie "vinceranno" — quelle che verranno scritte per ultime nella tabella hash. Supponiamo di aver chiamato il codice di inserimento con un array di coppie A/0 B/1 A/2 C/3 A/4. Quando il codice termina, le coppie B/1 e C/3 saranno garantite nella tabella, ma ci sarà anche una qualsiasi delle coppie A/0, A/2 o A/4. Questo può essere un problema oppure no: tutto dipende dall'applicazione. Puoi sapere in anticipo che nell'array di input non ci sono chiavi duplicate, oppure potrebbe non importarti quale valore è stato registrato per ultimo.
Se per te questo è un problema, allora dovresti separare le coppie duplicate in diverse chiamate al kernel CUDA. In CUDA, qualsiasi operazione con una chiamata al kernel termina sempre prima della successiva chiamata al kernel (almeno all'interno di un singolo thread. Nei thread diversi, i kernel vengono eseguiti in parallelo). Se nell'esempio sopra chiami un kernel con A/0 B/1 A/2 C/3, e un altro con A/4, allora la chiave A riceverà il valore 4.
Ora parliamo di se le funzioni lookup() e delete() debbano utilizzare un puntatore semplice (plain) o un puntatore variabile (volatile) per l'array di coppie nella hash table. afferma che:
Il compilatore può ottimizzare a sua discrezione le operazioni di lettura e scrittura nella memoria globale o condivisa... Queste ottimizzazioni possono essere disattivate utilizzando la parola chiave
volatile: ... qualsiasi riferimento a questa variabile viene compilato in un'istruzione reale di lettura o scrittura in memoria.
Considerazioni di correttezza non richiedono l'uso di volatile. Se il thread in esecuzione utilizza un valore memorizzato dalla precedente operazione di lettura, ciò significa che utilizzerà alcune informazioni leggermente obsolete. Tuttavia, queste informazioni provengono da uno stato valido della hash table in un dato momento della chiamata al kernel. Se hai bisogno di utilizzare le informazioni più aggiornate, puoi usare un puntatore volatile, ma la prestazione potrebbe diminuire leggermente: secondo i miei test, con 32 milioni di elementi da eliminare, la velocità è scesa da 500 milioni di eliminazioni al secondo a 450 milioni di eliminazioni al secondo.
Prestazioni
Nel test di inserimento di 64 milioni di elementi e nell'eliminazione di 32 milioni di essi, la competizione tra std::unordered_map e la hash table per GPU è praticamente assente:
std::unordered_map ha impiegato 70.691 ms per inserire ed eliminare elementi con successivo rilascio di unordered_map (il rilascio di milioni di elementi richiede tempo, poiché all'interno vengono eseguite numerose allocazioni di memoria). A dirla francamente, con unordered_map std:unordered_map std::unordered_map limitazioni completamente diverse. Questo è un unico flusso di esecuzione CPU, supporta chiavi e valori di qualsiasi dimensione, funziona bene con alti tassi di utilizzo e mostra prestazioni stabili dopo numerose cancellazioni.
La durata di funzionamento della tabella hash per GPU e interazione interprogrammatico è stata di 984 ms. Questo include il tempo impiegato per allocare la tabella in memoria e rimuoverla (un'unica allocazione di 1 GB di memoria, che in CUDA richiede del tempo), inserimento e cancellazione degli elementi, nonché iterazione su di essi. Sono stati considerati anche tutte le copie nella memoria e dalla memoria della scheda video.
Il funzionamento della tabella hash stessa ha impegnato 271 ms. Questo include il tempo speso dalla scheda video per inserire e rimuovere elementi e non considera il tempo per copiare in memoria e iterare sulla tabella risultante. Se la tabella GPU rimane attiva a lungo, o se la tabella hash è completamente contenuta nella memoria della scheda video (ad esempio, per creare una tabella hash che sarà utilizzata da altro codice GPU, non dalla CPU), il risultato del test è pertinente.
La tabella hash per la scheda video dimostra elevate prestazioni grazie all'elevata larghezza di banda e alla parallellizzazione attiva.
Svantaggi
L'architettura della tabella hash presenta diversi problemi di cui tenere conto:
- La clustering ostacola la sondatura lineare, il che porta a un posizionamento delle chiavi nella tabella non perfetto.
- Le chiavi non vengono rimosse utilizzando la funzione
deletee col tempo ingombrano la tabella.
Di conseguenza, le prestazioni della tabella hash possono diminuire gradualmente, specialmente se esiste a lungo e viene eseguito un numero elevato di inserimenti e cancellazioni. Uno dei modi per mitigare questi svantaggi è il rehashing in una nuova tabella con un coefficiente di utilizzo sufficientemente basso e filtro delle chiavi eliminate durante il rehashing.
Per illustrare i problemi descritti, utilizzo il codice sopra per creare una tabella di 128 milioni di elementi, inserendo ciclicamente 4 milioni di elementi finché non riempio 124 milioni di slot (coefficiente di utilizzo di circa 0,96). Ecco la tabella dei risultati, ogni riga rappresenta una chiamata al kernel CUDA con l'inserimento di 4 milioni di nuovi elementi in una tabella hash:
Coefficiente di utilizzo
Durata di inserimento di 4 194 304 elementi
0,00
11,608448 ms (361,314798 mln chiavi/seg.)
0,03
11,751424 ms (356,918799 mln chiavi/seg.)
0,06
11,942592 ms (351,205515 mln chiavi/seg.)
0,09
12,081120 ms (347,178429 mln chiavi/seg.)
0,12
12,242560 ms (342,600233 mln chiavi/seg.)
0,16
12,396448 ms (338,347235 mln chiavi/seg.)
0,19
12,533024 ms (334,660176 mln chiavi/seg.)
0,22
12,703328 ms (330,173626 mln chiavi/seg.)
0,25
12,884512 ms (325,530693 mln chiavi/seg.)
0,28
13,033472 ms (321,810182 mln chiavi/seg.)
0,31
13,239296 ms (316,807174 mln chiavi/seg.)
0,34
13,392448 ms (313,184256 mln chiavi/seg.)
0,37
13,624000 ms (307,861434 mln chiavi/seg.)
0,41
13,875520 ms (302,280855 mln chiavi/seg.)
0,44
14,126528 ms (296,909756 mln chiavi/seg.)
0,47
14,399328 ms (291,284699 mln chiavi/seg.)
0,50
14,690304 ms (285,515123 mln chiavi/seg.)
0,53
15,039136 ms (278,892623 mln chiavi/seg.)
0,56
15,478656 ms (270,973402 mln chiavi/seg.)
0,59
15,985664 ms (262,379092 mln chiavi/seg.)
0,62
16,668673 ms (251,627968 mln chiavi/seg.)
0,66
17,587200 ms (238,486174 mln chiavi/seg.)
0,69
18,690048 ms (224,413765 mln chiavi/seg.)
0,72
20,278816 ms (206,831789 mln chiavi/seg.)
0,75
22,545408 ms (186,038058 mln chiavi/seg.)
0,78
26,053312 ms (160,989275 mln chiavi/seg.)
0,81
31,895008 ms (131,503463 mln chiavi/seg.)
0,84
42,103294 ms (99,619378 mln chiavi/seg.)
0,87
61,849056 ms (67,815164 mln chiavi/seg.)
0,90
105,695999 ms (39,682713 mln chiavi/seg.)
0,94
240,204636 ms (17,461378 mln chiavi/seg.)
Man mano che il tasso di utilizzo cresce, le prestazioni diminuiscono. Ciò è indesiderato nella maggior parte dei casi. Se l'applicazione inserisce elementi in una tabella e poi li scarta (ad esempio, durante il conteggio delle parole in un libro), ciò non è un problema. Ma se l'applicazione utilizza una tabella hash a lungo termine (ad esempio, in un editor grafico per memorizzare porzioni non vuote delle immagini, quando l'utente inserisce e rimuove frequentemente informazioni), questo comportamento può causare problemi.
Ho misurato la profondità di sondaggio della tabella hash dopo 64 milioni di inserimenti (tasso di utilizzo 0,5). La profondità media era di 0,4774, quindi la maggior parte delle chiavi si trovava o nei migliori slot possibili, o a uno slot dalla posizione migliore. La profondità massima di sondaggio era di 60.
Poi ho misurato la profondità di sondaggio in una tabella con 124 milioni di inserimenti (tasso di utilizzo 0,97). La profondità media era già di 10,1757, e la massima era 6474 (!!). Le prestazioni del sondaggio lineare calano significativamente con alti tassi di utilizzo.
È meglio mantenere un basso fattore di utilizzo per questa tabella hash. Tuttavia, ciò aumenta le prestazioni a spese della memoria. Fortunatamente, nel caso di chiavi e valori a 32 bit, questo può essere giustificato. Se nel caso sopra citato la tabella con 128 milioni di elementi mantiene un fattore di utilizzo di 0,25, potremo posizionare al suo interno non più di 32 milioni di elementi, mentre altri 96 milioni di slot andranno persi — 8 byte per ogni coppia, 768 MB di memoria persa.
Si noti che stiamo parlando della perdita di memoria della scheda video, che è una risorsa più preziosa rispetto alla memoria di sistema. Anche se la maggior parte delle schede grafiche desktop moderne che supportano CUDA hanno almeno 4 GB di memoria (al momento della scrittura, la NVIDIA 2080 Ti ha 11 GB), perdere tali volumi non sarà la scelta più saggia.
In seguito scriverò in dettaglio sulla creazione di tabelle hash per schede video che non hanno problemi di profondità di probing, oltre a metodi per riutilizzare gli slot rimossi.
Misurazione della profondità di probing
Per determinare la profondità di probing di una chiave, possiamo estrarre l'hash della chiave (il suo indice ideale nella tabella) dal suo indice tabellare effettivo:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
A causa della magia di due numeri binari nel codice complementare e del fatto che la capacità della tabella hash è una potenza di due, questo approccio funzionerà anche quando l'indice della chiave viene spostato all'inizio della tabella. Prendiamo una chiave che è hashata in 1, ma inserita nello slot 3. Quindi, per una tabella con capacità 4, otteniamo (3 - 1) & 3, che è equivalente a 2.
Conclusione
Se hai domande o commenti, scrivimi in o apri un nuovo tema in.
Questo codice è stato scritto ispirandosi a magnifici articoli:
In futuro continuerò a scrivere su implementazioni di tabelle hash per schede video e analizzerò le loro prestazioni. Ho in programma la concatenazione, l'hashing di Robin Hood e l'hashing della cuculo utilizzando operazioni atomiche in strutture dati adatte alle schede video.
Fonte: habr.com
