L'aghju publicatu nantu à Github .
Hè una semplice tavola hash GPU capace di processà centinaie di milioni di inserimenti per seconda. Nantu à u mo laptop NVIDIA GTX 1060, u codice inserisce 64 milioni di coppie chjave-valori generati aleatoriamente in circa 210 ms è elimina 32 milioni di coppie in circa 64 ms.
Vale à dì, a velocità nantu à un laptop hè di circa 300 milioni di inserimenti / sec è 500 milioni di eliminazioni / sec.
A tavula hè scritta in CUDA, ancu s'è a stessa tecnica pò esse appiicata à HLSL o GLSL. L'implementazione hà parechje limitazioni per assicurà un altu rendimentu nantu à una carta video:
- Solu chjavi 32-bit è i stessi valori sò trattati.
- A tavula hash hà una dimensione fissa.
- E sta dimensione deve esse uguali à dui à u putere.
Per i chjavi è i valori, avete bisognu di riservà un marcatore di delimitatore simplice (in u codice sopra questu hè 0xffffffff).
Tavola Hash senza serrature
A tavola hash usa l'indirizzu apertu cù , vale à dì, hè solu un array di coppie chjave-valore chì hè guardatu in memoria è hà un rendimentu di cache superiore. U listessu ùn si pò dì per a catena, chì implica a ricerca di un punteru in una lista ligata. Un hash table hè un array simplice chì almacena elementi KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
A dimensione di a tavula hè una putenza di dui, micca un numeru primu, perchè una struzzione rapida hè abbastanza per applicà a maschera pow2 / AND, ma l'operatore di modulu hè assai più lento. Questu hè impurtante in u casu di sonda lineale, postu chì in una ricerca di tavula lineare l'indexu di slot deve esse impannillatu in ogni slot. È in u risultatu, u costu di l'operazione hè aghjuntu modulo in ogni slot.
A tavula guarda solu a chjave è u valore per ogni elementu, micca un hash di a chjave. Siccomu a tavula guarda solu chjave 32-bit, l'hash hè calculatu assai rapidamente. U codice sopra usa l'hash Murmur3, chì eseguisce solu uni pochi di turni, XOR è multiplicazioni.
A tavola hash usa tecniche di prutezzione di bloccu chì sò indipendenti di l'ordine di memoria. Ancu s'è alcune operazioni di scrittura disturbanu l'ordine di altre tali operazioni, a tavola hash mantene sempre u statu currettu. Avemu da parlà di questu quì sottu. A tecnica funziona bè cù carte video chì correnu migliaia di fili simultaneamente.
I chjavi è i valori in a tavola hash sò inizializzati per sviutata.
U codice pò esse mudificatu per trattà ancu e chjave di 64-bit è i valori. E chjave necessitanu operazioni atomiche di lettura, scrittura è comparazione è scambia. È i valori necessitanu operazioni atomiche di lettura è scrittura. Fortunatamente, in CUDA, l'operazioni di lettura-scrittura per i valori di 32 è 64 bit sò atomichi sempre chì sò naturali allinati (vede sottu). ), è e carte video muderne supportanu operazioni di comparazione è scambia atomica di 64 bit. Di sicuru, quandu si move à 64 bits, u rendiment diminuirà ligeramente.
Statu di a tavola Hash
Ogni coppia chjave-valore in una tavola hash pò avè unu di quattru stati:
- Chjave è valore sò vioti. In questu statu, a tavola hash hè inizializzata.
- A chjave hè stata scritta, ma u valore ùn hè ancu scrittu. Se un altru filu hè attualmente in lettura di dati, allora torna viotu. Questu hè normale, a listessa cosa avissi accadutu se un altru filu di l'esekzione avia travagliatu un pocu prima, è parlemu di una struttura di dati cuncurrenti.
- Sia a chjave è u valore sò arregistrati.
- U valore hè dispunibule per altri fili di esicuzzioni, ma a chjave ùn hè micca ancu. Questu pò accade perchè u mudellu di prugrammazione CUDA hà un mudellu di memoria urdinatu. Questu hè normale; in ogni casu, a chjave hè sempre viota, ancu s'ellu u valore ùn hè più cusì.
Una sfumatura impurtante hè chì una volta chì a chjave hè stata scritta à u slot, ùn si move più - ancu s'è a chjave hè sguassata, parlemu quì sottu.
U codice di a tavula di hash funziona ancu cù mudelli di memoria urdinati chì l'ordine in quale a memoria hè letta è scritta hè scunnisciutu. Quandu guardemu l'inserimentu, a ricerca è l'eliminazione in una table hash, ricordate chì ogni coppia chjave-valore hè in unu di i quattru stati descritti sopra.
Inserisce in una tavola hash
A funzione CUDA chì inserisce coppie chjave-valore in una tavola hash hè cusì:
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 inserisce una chjave, u codice iterate à traversu l'array di table hash cuminciendu cù l'hash di a chjave inserita. Ogni slot in l'array realiza una operazione atomica di paragunà è scambià chì paraguna a chjave in quellu slot à viotu. Se si rileva una mancata coincidenza, a chjave in u slot hè aghjurnata cù a chjave inserita, è poi a chjave di slot originale hè tornata. Se sta chjave originale era viota o currisponde à a chjave inserita, allura u codice truvò un slot adattatu per l'inserimentu è inserisce u valore inseritu in u slot.
Sè in una chjama di kernel gpu_hashtable_insert() Ci sò parechje elementi cù a listessa chjave, allora qualsiasi di i so valori ponu esse scrittu à u slot chjave. Questu hè cunsideratu normale: una di e scritture chjave-valore durante a chjama hà da successu, ma postu chì tuttu questu succede in parallelu in parechji fili di esecuzione, ùn pudemu micca predichendu quale scrittura di memoria serà l'ultima.
Ricerca di a tavola Hash
Codice per a ricerca di e chjave:
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 truvà u valore di una chjave guardata in una tavula, iteremu à traversu l'array cuminciendu cù l'hash di a chjave chì circhemu. In ogni slot, cuntrollemu se a chjave hè quella chì cercamu, è se cusì, vultemu u so valore. Avemu ancu verificatu se a chjave hè viota, è s'ellu hè cusì, abbandunemu a ricerca.
Se ùn pudemu micca truvà a chjave, u codice torna un valore viotu.
Tutte queste operazioni di ricerca ponu esse realizate simultaneamente attraversu inserimenti è eliminazioni. Ogni paru in a tavula avarà unu di i quattru stati descritti sopra per u flussu.
Eliminazione in una tavola hash
Codice per sguassà e chjave:
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 chjave hè fatta in modu inusual: lasciamu a chjave in a tavula è marcate u so valore (micca a chjave stessu) cum'è viotu. Stu codice hè assai simili à lookup(), salvu chì quandu una partita si trova nantu à una chjave, rende u so valore viotu.
Cumu l'esitatu sopra, una volta chì una chjave hè scritta in un slot, ùn hè più spustata. Ancu quandu un elementu hè sguassatu da a tavula, a chjave ferma in u locu, u so valore diventa solu viotu. Questu significa chì ùn avemu micca bisognu di utilizà una operazione di scrittura atomica per u valore di slot, perchè ùn importa micca chì u valore attuale hè viotu o micca - serà sempre viotu.
Ridimensionà una tavola hash
Pudete cambià a dimensione di una tavula hash creendu una tavola più grande è inserisce elementi micca vacanti da a vechja tavola in questu. Ùn aghju micca implementatu sta funziunalità perchè vulia mantene u codice di mostra simplice. Inoltre, in i prugrammi CUDA, l'allocazione di memoria hè spessu fatta in u codice di l'ospitu in quantu à u kernel CUDA.
L'articulu descrive cumu mudificà una struttura di dati cusì protetta da serratura.
Cumpetitività
In i frammenti di codice di funzione sopra gpu_hashtable_insert(), _lookup() и _delete() processà una coppia chjave-valore à un tempu. È più bassu gpu_hashtable_insert(), _lookup() и _delete() processà una serie di coppie in parallelu, ogni coppia in un filu di esecuzione GPU separatu:
// 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);
}
}
A tavola di hash resistente à i blocchi supporta inserimenti, ricerche è eliminazioni simultanee. Perchè i pariglii chjave-valori sò sempre in unu di quattru stati è i chjavi ùn si movenu micca, a tavula guarantisci a correzione ancu quandu diversi tipi di operazioni sò usati simultaneamente.
In ogni casu, se processemu un batch of insertions and deletions in parallel, è se l'array input di coppie cuntene chjavi duplicati, allora ùn pudemu micca predichendu quale coppie "vinceranu" - seranu scritte in l'ultima tavola hash. Dicemu chì avemu chjamatu u codice d'inserzione cù un array input di coppie A/0 B/1 A/2 C/3 A/4. Quandu u codice cumpleta, coppie B/1 и C/3 sò guarantiti per esse prisente in a tavula, ma à u stessu tempu qualsiasi di e coppie appariscenu in questu A/0, A/2 o A/4. Questu pò esse o micca un prublema - tuttu dipende di l'applicazione. Pudete sapè in anticipu chì ùn ci sò micca chjavi duplicati in l'array di input, o pudete micca cura di quale valore hè statu scrittu l'ultimu.
Se questu hè un prublema per voi, allora avete bisognu di separà e coppie duplicate in diverse chjamate di u sistema CUDA. In CUDA, ogni operazione chì chjamà u kernel hè sempre cumpletu prima di a prossima chjama di u kernel (almenu in un filu. In diversi fili, i kernels sò eseguiti in parallelu). In l'esempiu sopra, se chjamate un kernel cun A/0 B/1 A/2 C/3, è l'altru cun A/4, poi a chjave A uttene u valore 4.
Avà parlemu di se e funzioni duveranu lookup() и delete() Aduprate un punteru chjaru o volatile à un array di coppie in a tavola hash. Dici chì:
U compilatore pò sceglie per ottimisà e letture è scrive à a memoria globale o sparta... Queste ottimisazioni ponu esse disattivate usendu a keyword.
volatile: ... ogni riferimentu à sta variàbile hè cumpilatu in una memoria reale di leghje o scrive struzzione.
Considerazioni di correttezza ùn necessitanu micca applicazione volatile. Se u filu di l'esekzione usa un valore in cache da una operazione di lettura precedente, allora utilizerà infurmazione ligeramente obsoleta. Ma sempre, questu hè infurmazione da u statu currettu di a tavola hash in un certu mumentu di a chjama di u kernel. Sè avete bisognu di utilizà l'ultime informazioni, pudete aduprà l'indici volatile, ma tandu u rendiment diminuirà ligeramente: secondu i mo testi, quandu sguassate 32 milioni di elementi, a vitezza diminuì da 500 milioni di eliminazioni / sec à 450 milioni di eliminazioni / sec.
Produttività
In a prova per inserisce 64 milioni di elementi è sguassà 32 milioni di elli, a cumpetizione trà std::unordered_map è ùn ci hè quasi nisuna tabella di hash per a GPU:
std::unordered_map spentu 70 ms inserisce è caccià elementi è poi libbirà unordered_map (sbarazzà di milioni di elementi piglia assai tempu, perchè dentru unordered_map sò fatti parechje allocazioni di memoria). Onestamente parlà, std:unordered_map restrizioni completamente diverse. Hè un unicu filu di esecuzione di CPU, supporta i valori chjave di ogni dimensione, funziona bè à tassi d'utilizazione elevati, è mostra un rendimentu stabile dopu più eliminazioni.
A durata di a tavola hash per a GPU è a cumunicazione inter-programma era 984 ms. Questu include u tempu passatu à mette a tavula in memoria è sguassate (allocanu 1 GB di memoria una volta, chì piglia un pocu di tempu in CUDA), inserisce è sguassate elementi, è iterendu sopra. Tutte e copie da è da a memoria di a carta video sò ancu cunsiderate.
A tavola hash stessu hà pigliatu 271 ms per compie. Questu include u tempu passatu da a carta video chì inserisce è sguassate elementi, è ùn piglia micca in contu u tempu passatu à copià in memoria è iterendu nantu à a tavola resultanti. Se a tavola GPU vive per un bellu pezzu, o se a tavola hash hè cuntenuta sanu sanu in a memoria di a carta video (per esempiu, per creà una tavola hash chì serà utilizata da altre codice GPU è micca u processatore cintrali), allora u risultatu di a prova hè pertinente.
A tavola di hash per una carta video mostra un altu rendiment per via di un altu throughput è di parallelizazione attiva.
shortcomings
L'architettura di a tavola hash hà uni pochi di prublemi per esse cunzignati:
- A sonda lineale hè ostacolata da u clustering, chì face chì i chjavi in a tavula sò posti menu di perfetta.
- I chjave ùn sò micca eliminati cù a funzione
deleteè cù u tempu sbulicà a tavula.
In u risultatu, u funziunamentu di una tavola di hash pò degrade gradualmente, soprattuttu s'ellu esiste per un bellu pezzu è hà numerosi inseriti è sguassati. Una manera di mitigà questi svantaghji hè di ritruvà in una nova tavula cù un tassu d'utilizazione abbastanza bassu è filtrà i chjavi eliminati durante u rehashing.
Per illustrà i prublemi descritti, aduprà u codice sopra per creà una tavula cù 128 milioni d'elementi è un ciclu di 4 milioni d'elementi finu à ch'e aghju pienu 124 milioni di slots (tassu d'utilizazione di circa 0,96). Eccu a tabella di risultati, ogni fila hè una chjama di u kernel CUDA per inserisce 4 milioni di elementi novi in una tavola hash:
Tasso di usu
Durata d'inserzione 4 elementi
0,00
11,608448 ms (361,314798 milioni di tasti/sec.)
0,03
11,751424 ms (356,918799 milioni di tasti/sec.)
0,06
11,942592 ms (351,205515 milioni di tasti/sec.)
0,09
12,081120 ms (347,178429 milioni di tasti/sec.)
0,12
12,242560 ms (342,600233 milioni di tasti/sec.)
0,16
12,396448 ms (338,347235 milioni di tasti/sec.)
0,19
12,533024 ms (334,660176 milioni di tasti/sec.)
0,22
12,703328 ms (330,173626 milioni di tasti/sec.)
0,25
12,884512 ms (325,530693 milioni di tasti/sec.)
0,28
13,033472 ms (321,810182 milioni di tasti/sec.)
0,31
13,239296 ms (316,807174 milioni di tasti/sec.)
0,34
13,392448 ms (313,184256 milioni di tasti/sec.)
0,37
13,624000 ms (307,861434 milioni di tasti/sec.)
0,41
13,875520 ms (302,280855 milioni di tasti/sec.)
0,44
14,126528 ms (296,909756 milioni di tasti/sec.)
0,47
14,399328 ms (291,284699 milioni di tasti/sec.)
0,50
14,690304 ms (285,515123 milioni di tasti/sec.)
0,53
15,039136 ms (278,892623 milioni di tasti/sec.)
0,56
15,478656 ms (270,973402 milioni di tasti/sec.)
0,59
15,985664 ms (262,379092 milioni di tasti/sec.)
0,62
16,668673 ms (251,627968 milioni di tasti/sec.)
0,66
17,587200 ms (238,486174 milioni di tasti/sec.)
0,69
18,690048 ms (224,413765 milioni di tasti/sec.)
0,72
20,278816 ms (206,831789 milioni di tasti/sec.)
0,75
22,545408 ms (186,038058 milioni di tasti/sec.)
0,78
26,053312 ms (160,989275 milioni di tasti/sec.)
0,81
31,895008 ms (131,503463 milioni di tasti/sec.)
0,84
42,103294 ms (99,619378 milioni di tasti/sec.)
0,87
61,849056 ms (67,815164 milioni di tasti/sec.)
0,90
105,695999 ms (39,682713 milioni di tasti/sec.)
0,94
240,204636 ms (17,461378 milioni di tasti/sec.)
Quandu l'utilizazione aumenta, u rendiment diminuisce. Questu ùn hè micca desideratu in a maiò parte di i casi. Se una applicazione inserisce elementi in una tavula è poi li scarta (per esempiu, quandu cuntà e parolle in un libru), allora questu ùn hè micca un prublema. Ma se l'applicazione usa una tavola di hash longa (per esempiu, in un editore graficu per almacenà e parte micca vacanti di l'imaghjini induve l'utilizatore spessu inserisce è sguassate infurmazioni), allora stu cumpurtamentu pò esse problematicu.
È hà misuratu a prufundità di a prova di a tavola di hash dopu à 64 milioni di inserti (fattore d'utilizazione 0,5). A prufundità media era 0,4774, cusì a maiò parte di e chjave eranu o in u megliu slot pussibule o un slot luntanu da a megliu pusizioni. A prufundità massima di sonu era 60.
Allora aghju misuratu a prufundità di prufundità nantu à una tavula cù 124 milioni di inserti (fattore d'utilizazione 0,97). A prufundità media era digià 10,1757, è u massimu - 6474 (!!). U rendiment di sensazione lineare cala significativamente à tassi d'utilizazione elevati.
Hè megliu mantene a rata d'utilizazione di sta tavola hash bassa. Ma poi aumentemu u rendiment à a spesa di u cunsumu di memoria. Fortunatamente, in u casu di chjavi è valori 32-bit, questu pò esse ghjustificatu. Se in l'esempiu di sopra, in una tavula cù 128 milioni di elementi, mantenemu u fattore d'utilizazione di 0,25, allora ùn pudemu micca mette più di 32 milioni d'elementi in questu, è i 96 milioni di slot restante seranu persi - 8 bytes per ogni paru. , 768 MB di memoria persa.
Per piacè nutate chì parlemu di a perdita di memoria di a carta video, chì hè una risorsa più preziosa chè a memoria di u sistema. Ancu s'è a maiò parte di e carte grafiche di desktop moderni chì supportanu CUDA anu almenu 4 GB di memoria (à u mumentu di a scrittura, a NVIDIA 2080 Ti hà 11 GB), ùn saria micca sempre a decisione più sàvia di perde tali quantità.
In seguitu scriveraghju più nantu à a creazione di tavule di hash per e carte video chì ùn anu micca prublemi cù a prufundità di sonda, è ancu modi per riutilizà slots eliminati.
Misura di a prufundità di u sonu
Per determinà a prufundità di prufundità di una chjave, pudemu estrarre l'hash di a chjave (u so indice di tavula ideale) da u so indice di tabella attuale:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
A causa di a magia di dui numeri binari cumplementarii di dui è u fattu chì a capacità di a tavula hash hè duie à u putere di dui, stu approcciu hà da travaglià ancu quandu l'indici chjave hè spustatu à u principiu di a tavola. Pigliemu una chjave chì hà chjapputu à 1, ma hè inseritu in u slot 3. Allora per una tavola cù capacità 4 avemu avutu. (3 — 1) & 3, chì hè equivalente à 2.
cunchiusioni
Sì avete dumande o cumenti, per piacè email à mè à o apre un novu tema in.
Stu codice hè statu scrittu sottu ispirazioni da articuli eccellenti:
In u futuru, continueraghju à scrive nantu à implementazioni di tavule di hash per carte video è analizà a so prestazione. I mo piani includenu a catena, l'hashing Robin Hood, è l'hashing cuckoo utilizendu operazioni atomiche in strutture di dati chì sò GPU amichevuli.
Source: www.habr.com