L'ho pubblicato su Github .
Si tratta di una semplice tabella hash GPU in grado di elaborare centinaia di milioni di inserimenti al secondo. Sul mio laptop 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.
Cioè, la velocità su un laptop è di circa 300 milioni di inserimenti/sec e 500 milioni di eliminazioni/sec.
La tabella è scritta in CUDA, sebbene la stessa tecnica possa essere applicata a HLSL o GLSL. L'implementazione presenta diverse limitazioni per garantire prestazioni elevate su una scheda video:
- Vengono elaborate solo chiavi a 32 bit e gli stessi valori.
- La tabella hash ha una dimensione fissa.
- E questa dimensione deve essere pari a due alla potenza.
Per chiavi e valori, è necessario riservare un semplice indicatore delimitatore (nel codice sopra è 0xffffffff).
Tabella hash senza blocchi
La tabella hash utilizza l'indirizzamento aperto con , ovvero si tratta semplicemente di un array di coppie chiave-valore archiviato in memoria e con prestazioni della cache superiori. Lo stesso non si può dire per il concatenamento, che implica la ricerca di un puntatore in un elenco collegato. Una tabella hash è un semplice array che memorizza elementi KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
La dimensione della tabella è una potenza di due, non un numero primo, perché un'istruzione veloce è sufficiente per applicare la maschera pow2/AND, ma l'operatore modulo è molto più lento. Ciò è importante nel caso del sondaggio lineare, poiché in una ricerca su tabella lineare l'indice dello slot deve essere racchiuso in ogni slot. Di conseguenza, il costo dell'operazione viene aggiunto modulo in ogni slot.
La tabella memorizza solo la chiave e il valore per ciascun elemento, non un hash della chiave. Poiché la tabella memorizza solo chiavi a 32 bit, l'hash viene calcolato molto rapidamente. Il codice sopra utilizza l'hash Murmur3, che esegue solo pochi spostamenti, XOR e moltiplicazioni.
La tabella hash utilizza tecniche di protezione con blocco indipendenti dall'ordine della memoria. Anche se alcune operazioni di scrittura interrompono l'ordine di altre operazioni simili, la tabella hash manterrà comunque lo stato corretto. Ne parleremo di seguito. La tecnica funziona alla grande con le schede video che eseguono migliaia di thread contemporaneamente.
Le chiavi e i valori nella tabella hash vengono inizializzati su vuoti.
Il codice può essere modificato per gestire anche chiavi e valori a 64 bit. Le chiavi richiedono operazioni atomiche di lettura, scrittura e confronto e scambio. E i valori richiedono operazioni di lettura e scrittura atomiche. Fortunatamente, in CUDA, le operazioni di lettura-scrittura per valori a 32 e 64 bit sono atomiche purché siano allineate naturalmente (vedi sotto). ) e le moderne schede video supportano operazioni di confronto e scambio atomico a 64 bit. Naturalmente, passando a 64 bit, le prestazioni diminuiranno leggermente.
Stato della tabella hash
Ciascuna coppia chiave-valore in una tabella hash può avere uno dei quattro stati:
- Chiave e valore sono vuoti. In questo stato, la tabella hash viene inizializzata.
- La chiave è stata annotata, ma il valore non è stato ancora scritto. Se un altro thread sta attualmente leggendo dati, restituisce vuoto. Questo è normale, la stessa cosa sarebbe accaduta se un altro thread di esecuzione avesse funzionato un po' prima, e stiamo parlando di una struttura dati concorrente.
- Vengono registrati sia la chiave che il valore.
- Il valore è disponibile per altri thread di esecuzione, ma la chiave non lo è ancora. Ciò può accadere perché il modello di programmazione CUDA ha un modello di memoria poco ordinato. Questo è normale, in ogni caso la chiave è ancora vuota anche se il valore non lo è più.
Una sfumatura importante è che una volta che la chiave è stata scritta nello slot, non si muove più - anche se la chiave viene cancellata, ne parleremo di seguito.
Il codice della tabella hash funziona anche con modelli di memoria poco ordinati in cui l'ordine in cui la memoria viene letta e scritta è sconosciuto. Mentre osserviamo l'inserimento, la ricerca e l'eliminazione in una tabella hash, ricorda che ciascuna coppia chiave-valore si trova in uno dei quattro stati descritti sopra.
Inserimento in una tabella hash
La funzione CUDA che inserisce coppie chiave-valore in una tabella hash si presenta 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 scorre l'array della tabella hash iniziando con l'hash della chiave inserita. Ogni slot nell'array esegue un'operazione atomica di confronto e scambio che confronta la chiave in quello slot con quella vuota. Se viene rilevata una mancata corrispondenza, la chiave nello slot viene aggiornata con la chiave inserita, quindi viene restituita la chiave dello slot originale. Se questa chiave originale era vuota o corrispondeva alla chiave inserita, il codice ha trovato uno slot adatto per l'inserimento e ha inserito il valore inserito nello slot.
Se in una chiamata del kernel gpu_hashtable_insert() ci sono più elementi con la stessa chiave, quindi qualsiasi loro valore può essere scritto nello slot della chiave. Questo è considerato normale: una delle scritture di valori-chiave durante la chiamata avrà successo, ma poiché tutto ciò avviene in parallelo all'interno di più thread di esecuzione, non possiamo prevedere quale 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 in una tabella, iteriamo attraverso l'array iniziando con l'hash della chiave che stiamo cercando. In ogni slot controlliamo se la chiave è quella che stiamo cercando e, in tal caso, restituiamo il suo valore. Controlliamo anche se la chiave è vuota e, in tal caso, interrompiamo la ricerca.
Se non riusciamo a trovare la chiave, il codice restituisce un valore vuoto.
Tutte queste operazioni di ricerca possono essere eseguite contemporaneamente tramite inserimenti e cancellazioni. Ciascuna coppia nella tabella avrà uno dei quattro stati sopra descritti per il flusso.
Eliminazione in una tabella hash
Codice per eliminare le 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);
}
}
L'eliminazione di una chiave avviene in un 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 su una chiave, il suo valore viene vuoto.
Come accennato in precedenza, una volta scritta una chiave in uno slot, non viene più spostata. Anche quando un elemento viene eliminato dalla tabella, la chiave rimane al suo posto, il suo valore diventa semplicemente vuoto. Ciò significa che non è necessario utilizzare un'operazione di scrittura atomica per il valore dello slot, perché non importa se il valore corrente è vuoto o meno: lo diventerà comunque.
Ridimensionamento di una tabella hash
È possibile modificare la dimensione di una tabella hash creando una tabella più grande e inserendovi elementi non vuoti della vecchia tabella. Non ho implementato questa funzionalità perché volevo mantenere semplice il codice di esempio. Inoltre, nei programmi CUDA, l'allocazione della memoria viene spesso effettuata nel codice host anziché nel kernel CUDA.
l'articolo descrive come modificare tale struttura dati protetta da blocco.
Competitività
Negli snippet di codice della funzione sopra gpu_hashtable_insert(), _lookup() и _delete() elaborare una coppia chiave-valore alla volta. E più in basso gpu_hashtable_insert(), _lookup() и _delete() elabora una serie di coppie in parallelo, ciascuna 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);
}
}
La tabella hash resistente ai blocchi supporta inserimenti, ricerche ed eliminazioni simultanee. 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 quando vengono utilizzati contemporaneamente diversi tipi di operazioni.
Tuttavia, se elaboriamo un batch di inserimenti ed eliminazioni in parallelo e se l'array di coppie di input contiene chiavi duplicate, non saremo in grado di prevedere quali coppie "vinceranno": verranno scritte per ultime nella tabella hash. Diciamo che abbiamo chiamato il codice di inserimento con un array di coppie di input A/0 B/1 A/2 C/3 A/4. Una volta completato il codice, si accoppia B/1 и C/3 sono garantiti nella tabella, ma allo stesso tempo in essa apparirà una qualsiasi delle coppie A/0, A/2 o A/4. Questo può essere o meno un problema: tutto dipende dall'applicazione. Potresti sapere in anticipo che non ci sono chiavi duplicate nell'array di input o potresti non interessarti quale valore è stato scritto per ultimo.
Se questo è un problema per te, devi separare le coppie duplicate in diverse chiamate di sistema CUDA. In CUDA, qualsiasi operazione che chiama il kernel viene sempre completata prima della successiva chiamata al kernel (almeno all'interno di un thread. In thread diversi, i kernel vengono eseguiti in parallelo). Nell'esempio sopra, se chiami un kernel con A/0 B/1 A/2 C/3, e l'altro con A/4, quindi la chiave A otterrà il valore 4.
Ora parliamo se le funzioni dovrebbero lookup() и delete() utilizzare un puntatore semplice o volatile a un array di coppie nella tabella hash. Afferma che:
Il compilatore può scegliere di ottimizzare le letture e le scritture sulla memoria globale o condivisa... Queste ottimizzazioni possono essere disabilitate utilizzando la parola chiave
volatile: ... qualsiasi riferimento a questa variabile viene compilato in un'istruzione di lettura o scrittura della memoria reale.
Le considerazioni sulla correttezza non richiedono applicazione volatile. Se il thread di esecuzione utilizza un valore memorizzato nella cache da un'operazione di lettura precedente, utilizzerà informazioni leggermente obsolete. Ma si tratta comunque di informazioni dallo stato corretto della tabella hash in un determinato momento della chiamata al kernel. Se è necessario utilizzare le informazioni più recenti, è possibile utilizzare l'indice volatile, ma poi le prestazioni diminuiranno leggermente: secondo i miei test, eliminando 32 milioni di elementi, la velocità è scesa da 500 milioni di eliminazioni/sec a 450 milioni di eliminazioni/sec.
Производительность
Nel test per l'inserimento di 64 milioni di elementi e l'eliminazione di 32 milioni di essi, competizione tra std::unordered_map e praticamente non esiste una tabella hash per la GPU:
std::unordered_map ha impiegato 70 ms inserendo e rimuovendo elementi e quindi liberandoli unordered_map (sbarazzarsi di milioni di elementi richiede molto tempo, perché inside unordered_map vengono effettuate più allocazioni di memoria). Francamente, std:unordered_map restrizioni completamente diverse. È un singolo thread di esecuzione della CPU, supporta valori-chiave di qualsiasi dimensione, funziona bene con tassi di utilizzo elevati e mostra prestazioni stabili dopo più eliminazioni.
La durata della tabella hash per la GPU e la comunicazione tra programmi è stata di 984 ms. Ciò include il tempo impiegato per inserire la tabella in memoria ed eliminarla (allocando 1 GB di memoria una volta, il che richiede un po' di tempo in CUDA), inserendo ed eliminando elementi e ripetendoli. Vengono prese in considerazione anche tutte le copie da e verso la memoria della scheda video.
Il completamento della tabella hash stessa ha richiesto 271 ms. Ciò include il tempo impiegato dalla scheda video per inserire ed eliminare elementi e non tiene conto del tempo impiegato per copiare in memoria ed eseguire iterazioni sulla tabella risultante. Se la tabella GPU dura a lungo, o se la tabella hash è contenuta interamente nella memoria della scheda video (ad esempio, per creare una tabella hash che verrà utilizzata da altro codice GPU e non dal processore centrale), allora il risultato del test è rilevante.
La tabella hash per una scheda video dimostra prestazioni elevate grazie all'elevato throughput e alla parallelizzazione attiva.
Limitazioni
L'architettura della tabella hash presenta alcuni problemi di cui tenere conto:
- Il sondaggio lineare è ostacolato dal clustering, che fa sì che le chiavi nella tabella non siano posizionate perfettamente.
- Le chiavi non vengono rimosse utilizzando la funzione
deletee col tempo ingombrano la tavola.
Di conseguenza, le prestazioni di una tabella hash possono gradualmente peggiorare, soprattutto se esiste da molto tempo e presenta numerosi inserimenti ed eliminazioni. Un modo per mitigare questi svantaggi è ripetere l'hashing in una nuova tabella con un tasso di utilizzo abbastanza basso e filtrare le chiavi rimosse durante il rehashing.
Per illustrare i problemi descritti, utilizzerò il codice precedente per creare una tabella con 128 milioni di elementi e scorrere 4 milioni di elementi fino a riempire 124 milioni di slot (tasso di utilizzo di circa 0,96). Ecco la tabella dei risultati, ogni riga è una chiamata del kernel CUDA per inserire 4 milioni di nuovi elementi in una tabella hash:
Fattore di utilizzo
Durata inserimento 4 elementi
0,00
11,608448 ms (361,314798 milioni di chiavi/sec.)
0,03
11,751424 ms (356,918799 milioni di chiavi/sec.)
0,06
11,942592 ms (351,205515 milioni di chiavi/sec.)
0,09
12,081120 ms (347,178429 milioni di chiavi/sec.)
0,12
12,242560 ms (342,600233 milioni di chiavi/sec.)
0,16
12,396448 ms (338,347235 milioni di chiavi/sec.)
0,19
12,533024 ms (334,660176 milioni di chiavi/sec.)
0,22
12,703328 ms (330,173626 milioni di chiavi/sec.)
0,25
12,884512 ms (325,530693 milioni di chiavi/sec.)
0,28
13,033472 ms (321,810182 milioni di chiavi/sec.)
0,31
13,239296 ms (316,807174 milioni di chiavi/sec.)
0,34
13,392448 ms (313,184256 milioni di chiavi/sec.)
0,37
13,624000 ms (307,861434 milioni di chiavi/sec.)
0,41
13,875520 ms (302,280855 milioni di chiavi/sec.)
0,44
14,126528 ms (296,909756 milioni di chiavi/sec.)
0,47
14,399328 ms (291,284699 milioni di chiavi/sec.)
0,50
14,690304 ms (285,515123 milioni di chiavi/sec.)
0,53
15,039136 ms (278,892623 milioni di chiavi/sec.)
0,56
15,478656 ms (270,973402 milioni di chiavi/sec.)
0,59
15,985664 ms (262,379092 milioni di chiavi/sec.)
0,62
16,668673 ms (251,627968 milioni di chiavi/sec.)
0,66
17,587200 ms (238,486174 milioni di chiavi/sec.)
0,69
18,690048 ms (224,413765 milioni di chiavi/sec.)
0,72
20,278816 ms (206,831789 milioni di chiavi/sec.)
0,75
22,545408 ms (186,038058 milioni di chiavi/sec.)
0,78
26,053312 ms (160,989275 milioni di chiavi/sec.)
0,81
31,895008 ms (131,503463 milioni di chiavi/sec.)
0,84
42,103294 ms (99,619378 milioni di chiavi/sec.)
0,87
61,849056 ms (67,815164 milioni di chiavi/sec.)
0,90
105,695999 ms (39,682713 milioni di chiavi/sec.)
0,94
240,204636 ms (17,461378 milioni di chiavi/sec.)
All’aumentare dell’utilizzo, le prestazioni diminuiscono. Ciò non è auspicabile nella maggior parte dei casi. Se un'applicazione inserisce elementi in una tabella e poi li scarta (ad esempio, quando conta le parole in un libro), questo non è un problema. Ma se l'applicazione utilizza una tabella hash di lunga durata (ad esempio, in un editor grafico per memorizzare parti non vuote di immagini in cui l'utente inserisce ed elimina frequentemente informazioni), allora questo comportamento può essere problematico.
E ha misurato la profondità di sondaggio della tabella hash dopo 64 milioni di inserimenti (fattore di utilizzo 0,5). La profondità media era 0,4774, quindi la maggior parte dei tasti si trovava nel miglior slot possibile o a uno slot di distanza dalla posizione migliore. La profondità massima di sondaggio era 60.
Ho poi misurato la profondità di sondaggio su una tavola con 124 milioni di inserti (fattore di utilizzo 0,97). La profondità media era già 10,1757, e la massima - 6474 (!!). Le prestazioni del rilevamento lineare diminuiscono significativamente con tassi di utilizzo elevati.
È meglio mantenere basso il tasso di utilizzo di questa tabella hash. Ma poi aumentiamo le prestazioni a scapito del consumo di memoria. Fortunatamente, nel caso di chiavi e valori a 32 bit, ciò può essere giustificato. Se nell'esempio sopra, in una tabella con 128 milioni di elementi, manteniamo il fattore di utilizzo di 0,25, non possiamo inserire più di 32 milioni di elementi e i restanti 96 milioni di slot andranno persi - 8 byte per ogni coppia , 768 MB di memoria persa.
Tieni presente che stiamo parlando della perdita di memoria della scheda video, che è una risorsa più preziosa della memoria di sistema. Anche se la maggior parte delle moderne schede grafiche desktop che supportano CUDA hanno almeno 4 GB di memoria (al momento in cui scriviamo, la NVIDIA 2080 Ti ha 11 GB), non sarebbe comunque la decisione più saggia perdere tali quantità.
Successivamente scriverò di più sulla creazione di tabelle hash per schede video che non presentano problemi con la profondità di sondaggio, nonché sui modi per riutilizzare gli slot eliminati.
Misurazione della profondità del sondaggio
Per determinare la profondità di sondaggio di una chiave, possiamo estrarre l'hash della chiave (il suo indice di tabella ideale) dal suo indice di tabella 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 dei numeri binari in complemento a due e del fatto che la capacità della tabella hash è due alla potenza di due, questo approccio funzionerà anche quando l'indice della chiave viene spostato all'inizio della tabella. Prendiamo una chiave con hash 1, ma inserita nello slot 3. Quindi per una tabella con capacità 4 otteniamo (3 — 1) & 3, che equivale a 2.
conclusione
Se avete domande o commenti, scrivetemi a o apri un nuovo argomento in.
Questo codice è stato scritto ispirandosi ad articoli eccellenti:
In futuro continuerò a scrivere sulle implementazioni delle tabelle hash per le schede video e ad analizzarne le prestazioni. I miei piani includono il concatenamento, l'hashing di Robin Hood e l'hashing del cuculo utilizzando operazioni atomiche in strutture di dati compatibili con la GPU.
Fonte: habr.com