Ech hunn et op Github gepost .
Et ass en einfachen GPU Hash-Tabelle fäeg fir Honnerte vu Millioune Inserts pro Sekonn ze veraarbecht. Op mengem NVIDIA GTX 1060 Laptop setzt de Code 64 Millioune zoufälleg generéiert Schlësselwäertpaaren an ongeféier 210 ms a läscht 32 Millioune Pairen an ongeféier 64 ms.
Dat ass, d'Geschwindegkeet op engem Laptop ass ongeféier 300 Milliounen Inserts / Sek an 500 Milliounen Läschen / Sek.
Den Dësch ass an CUDA geschriwwen, obwuel déi selwecht Technik kann op HLSL oder GLSL applizéiert ginn. D'Implementatioun huet verschidde Aschränkungen fir eng héich Leeschtung op enger Videokaart ze garantéieren:
- Nëmmen 32-Bit Schlësselen an déiselwecht Wäerter ginn veraarbecht.
- Den Hash Dësch huet eng fix Gréisst.
- An dës Gréisst muss d'selwecht sinn wéi zwee fir d'Muecht.
Fir Schlësselen a Wäerter musst Dir en einfachen Ofgrenzungsmarker reservéieren (am uewe genannte Code ass dëst 0xffffffff).
Hash Dësch ouni Schleisen
Den Hash-Tabelle benotzt oppen Adresséierung mat , dat heescht, et ass einfach eng Rei vu Schlësselwäertpaaren déi an der Erënnerung gespäichert sinn an eng super Cache Leeschtung huet. Datselwecht kann net fir Ketten gesot ginn, wat involvéiert d'Sich no engem Pointer an enger verlinkter Lëscht. En Hash Dësch ass eng einfach Array déi Elementer späichert KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
D'Gréisst vum Dësch ass eng Kraaft vun zwee, net eng Primzuel, well eng séier Instruktioun ass genuch fir d'pow2 / AND Mask z'applizéieren, awer de Modulusoperateur ass vill méi lues. Dëst ass wichteg am Fall vu linearer Sonde, well an enger linearer Tabelle Lookup muss de Slotindex an all Slot gewéckelt ginn. An als Resultat ginn d'Käschte vun der Operatioun modulo an all Slot bäigefüügt.
Den Dësch späichert nëmmen de Schlëssel a Wäert fir all Element, net en Hash vum Schlëssel. Well den Dësch nëmmen 32-Bit Schlësselen späichert, gëtt den Hash ganz séier berechent. De Code uewen benotzt de Murmur3 Hash, deen nëmmen e puer Verréckelungen, XORs a Multiplikatioune mécht.
Den Hash-Tabelle benotzt Sperrschutztechniken déi onofhängeg vun der Erënnerungsuerdnung sinn. Och wann e puer Schreifoperatioune d'Uerdnung vun aneren esou Operatiounen stéieren, wäert d'Hash-Table nach ëmmer de richtege Staat behalen. Mir schwätzen iwwer dëst ënnert. D'Technik funktionnéiert gutt mat Videokaarten déi Dausende vu Threads gläichzäiteg lafen.
D'Schlësselen a Wäerter an der Hash-Tabelle ginn initialiséiert fir eidel.
De Code kann geännert ginn fir och 64-Bit Schlësselen a Wäerter ze handhaben. Schlësselen erfuerderen atomesch Liesen, Schreiwen a Vergläichen-a-Swap Operatiounen. A Wäerter erfuerderen atomarer Lies- a Schreifoperatioune. Glécklecherweis, am CUDA, liesen-schreiwen Operatioune fir 32- a 64-Bit Wäerter sinn atomar soulaang se natierlech ausgeriicht sinn (kuckt hei ënnen). ), a modern Videokaarten ënnerstëtzen 64-Bit Atomvergläich-an-Austausch Operatiounen. Natierlech, wann Dir op 64 Bits bewegt, wäert d'Performance liicht erofgoen.
Hash Dësch Staat
All Schlësselwäertpaar an enger Hash-Tabelle kann ee vu véier Staaten hunn:
- Schlëssel a Wäert sinn eidel. An dësem Zoustand gëtt den Hash-Tabelle initialiséiert.
- De Schlëssel ass opgeschriwwen, awer de Wäert ass nach net geschriwwe ginn. Wann en anere Fuedem am Moment Daten liest, gëtt se dann eidel zréck. Dëst ass normal, datselwecht wier geschitt wann en anere Fuedem vun der Ausféierung e bësse méi fréi geschafft hätt, a mir schwätzen iwwer eng concurrent Datestruktur.
- Souwuel de Schlëssel wéi och de Wäert ginn opgeholl.
- De Wäert ass verfügbar fir aner Ausféierungsfäegkeeten, awer de Schlëssel ass nach net. Dëst ka geschéien well de CUDA Programméierungsmodell e locker bestallt Erënnerungsmodell huet. Dëst ass normal; op alle Fall ass de Schlëssel nach ëmmer eidel, och wann de Wäert net méi esou ass.
Eng wichteg Nuance ass datt eemol de Schlëssel op de Slot geschriwwe gouf, et net méi bewegt - och wann de Schlëssel geläscht gëtt, wäerte mir hei drënner schwätzen.
Den Hash Table Code funktionnéiert souguer mat locker bestallten Erënnerungsmodeller an deenen d'Uerdnung an där d'Erënnerung gelies a geschriwwe gëtt onbekannt ass. Wéi mir d'Insertion, Lookup an d'Läschen an enger Hash-Tabel kucken, erënnert datt all Schlësselwäertpaar an engem vun de véier uewen beschriwwene Staaten ass.
Asetzen an eng Hash-Tabelle
D'CUDA Funktioun déi Schlëssel-Wäertpaaren an eng Hash-Tabel setzt gesäit esou aus:
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);
}
}
Fir e Schlëssel anzeginn, iteréiert de Code duerch d'Hash-Tabelle-Array unzefänken mam Hash vum agebaute Schlëssel. All Slot am Array mécht eng atomar Compare-and-Swap Operatioun déi de Schlëssel an deem Slot vergläicht fir eidel. Wann e Mëssverständnis festgestallt gëtt, gëtt de Schlëssel am Slot mat dem agebaute Schlëssel aktualiséiert, an dann gëtt den urspréngleche Slot Schlëssel zréck. Wann dësen urspréngleche Schlëssel eidel war oder mat dem agebaute Schlëssel entsprécht, dann huet de Code e passende Slot fir d'Insertioun fonnt an den agebaute Wäert an de Slot agebaut.
Wann an engem Kärel Opruff gpu_hashtable_insert() et gi verschidde Elementer mam selwechte Schlëssel, da kann ee vun hire Wäerter op de Schlësselschlëssel geschriwwe ginn. Dëst gëtt als normal ugesinn: ee vun de Schlësselwäert Schreiwen wärend dem Uruff geléngt, awer well dat alles parallel an e puer Ausféierungsfäegkeeten geschitt, kënne mir net viraussoen wéi eng Erënnerungsschrëft déi lescht wäert sinn.
Hash Dësch Lookup
Code fir Schlësselen ze sichen:
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);
}
}
Fir de Wäert vun engem Schlëssel ze fannen, deen an enger Tabell gelagert ass, iteréiere mir duerch d'Array unzefänken mam Hash vum Schlëssel dee mir sichen. An all Slot kontrolléieren mir ob de Schlëssel deen ass deen mir sichen, a wa jo, mir ginn säi Wäert zréck. Mir kontrolléieren och ob de Schlëssel eidel ass, a wa jo, ofbriechen mir d'Sich.
Wa mir de Schlëssel net fannen, gëtt de Code en eidele Wäert zréck.
All dës Sichoperatioune kënnen gläichzäiteg duerch Insertiounen a Läschen duerchgefouert ginn. All Pair an der Tabell wäert ee vun de véier Staaten uewen beschriwwen fir de Flux hunn.
Läschen an engem Hash Dësch
Code fir Schlësselen ze läschen:
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);
}
}
E Schlëssel läschen gëtt op eng ongewéinlech Manéier gemaach: mir loossen de Schlëssel an der Tabell a markéieren säi Wäert (net de Schlëssel selwer) als eidel. Dëse Code ass ganz ähnlech wéi lookup(), ausser datt wann e Match op engem Schlëssel fonnt gëtt, mécht säi Wäert eidel.
Wéi uewen ernimmt, eemol e Schlëssel op e Slot geschriwwe gëtt, gëtt et net méi geplënnert. Och wann en Element vum Dësch geläscht gëtt, bleift de Schlëssel op der Plaz, säi Wäert gëtt einfach eidel. Dat heescht, datt mir brauchen net eng atomarer Schreiwen Operatioun fir Slot Wäert ze benotzen, well et egal ob den aktuelle Wäert eidel ass oder net - et wäert nach eidel ginn.
Änneren d'Gréisst vun engem Hash Dësch
Dir kënnt d'Gréisst vun engem Hash-Tabelle änneren andeems Dir e méi groussen Dësch erstellt an net eidel Elementer aus der aler Tabelle setzt. Ech hunn dës Funktionalitéit net ëmgesat well ech wollt de Probecode einfach halen. Ausserdeem, an CUDA Programmer, gëtt d'Erënnerungsallokatioun dacks am Hostcode gemaach anstatt am CUDA Kernel.
Den Artikel beschreift wéi esou eng gespaarten geschützt Date Struktur änneren.
Kompetitivitéit
An der uewen Funktioun Code Snippets gpu_hashtable_insert(), _lookup() и _delete() veraarbecht ee Schlëssel-Wäertpaar gläichzäiteg. An méi niddereg gpu_hashtable_insert(), _lookup() и _delete() veraarbecht eng Array vu Pairen parallel, all Paar an engem getrennten GPU Ausféierungs thread:
// 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);
}
}
De gespaarten-resistente Hash-Tabelle ënnerstëtzt gläichzäiteg Inserts, Lookups a Läschen. Well Schlëssel-Wäertpairen ëmmer an engem vu véier Staaten sinn an d'Schlësselen net beweegen, garantéiert den Dësch Richtegkeet och wann verschidden Aarte vu Operatiounen gläichzäiteg benotzt ginn.
Wéi och ëmmer, wa mir e Batch vun Insertiounen a Läschen parallel veraarbechten, a wann d'Input-Array vu Pairen duplizéiert Schlësselen enthält, da kënne mir net viraussoen wéi eng Pairen "gewannen" wäerte fir d'lescht op d'Hash-Tabelle geschriwwe ginn. Loosst eis soen datt mir den Insertiounscode mat enger Input-Array vu Pairen genannt hunn A/0 B/1 A/2 C/3 A/4. Wann de Code fäerdeg ass, Pairen B/1 и C/3 sinn garantéiert an der Tabell präsent ze sinn, awer gläichzäiteg wäert iergendeen vun de Puer dran erschéngen A/0, A/2 oder A/4. Dëst kann e Problem sinn oder net - et hänkt alles vun der Applikatioun of. Dir wësst vläicht am Viraus datt et keng Duplikatschlësselen an der Input-Array sinn, oder Dir kënnt egal wéi ee Wäert lescht geschriwwe gouf.
Wann dëst e Problem fir Iech ass, da musst Dir d'Duplikatpaar a verschidde CUDA System Uruff trennen. Am CUDA gëtt all Operatioun, déi de Kärel nennt, ëmmer virum nächste Kernelruff ofgeschloss (op d'mannst bannent engem Thread. A verschiddene Threads ginn d'Kären parallel ausgefouert). Am Beispill uewen, wann Dir e Kärel rufft mat A/0 B/1 A/2 C/3, an déi aner mat A/4, dann de Schlëssel A wäert de Wäert kréien 4.
Elo schwätze mer iwwer ob Funktiounen sollen lookup() и delete() benotzt en einfachen oder liichtflüchtege Pointer op eng Rei vu Pairen an der Hash-Tabelle. Staaten déi:
De Compiler kann wielen fir Liesen a Schreiwen op global oder gedeelt Erënnerung ze optimiséieren ... Dës Optimisatioune kënne mat dem Schlësselwuert deaktivéiert ginn
volatile: ... all Referenz op dës Variabel gëtt an eng richteg Erënnerung Lies- oder Schreifinstruktioun zesummegesat.
Richtegkeet Considératiounen verlaangen keng Applikatioun volatile. Wann den Ausféierungsfuedem e cachéierte Wäert vun enger fréierer Liesoperatioun benotzt, da wäert et liicht verännert Informatioun benotzen. Awer trotzdem ass dëst Informatioun aus dem korrekten Zoustand vun der Hash-Tabelle zu engem bestëmmte Moment vum Kernelruff. Wann Dir déi lescht Informatioun benotze musst, kënnt Dir den Index benotzen volatile, mee da wäert d'Performance liicht erofgoen: no mengen Tester, wann Dir 32 Milliounen Elementer geläscht hutt, ass d'Geschwindegkeet vu 500 Millioune Läschen / sec op 450 Millioune Läschen / sec erofgaang.
Produktivitéit
Am Test fir 64 Milliounen Elementer ze setzen an 32 Millioune vun hinnen ze läschen, Konkurrenz tëscht std::unordered_map an et gëtt quasi keen Hash Dësch fir d'GPU:
std::unordered_map 70 ms verbrauchen Elementer ze setzen an ze läschen an se dann ze befreien unordered_map (Millioune vun Elementer entgoen hëlt vill Zäit, well dobannen unordered_map Multiple Erënnerungsallokatiounen ginn gemaach). Éierlech gesot, std:unordered_map komplett aner Restriktiounen. Et ass en eenzegen CPU Fuedem vun der Ausféierung, ënnerstëtzt Schlësselwäerter vun all Gréisst, funktionnéiert gutt bei héijen Notzungsraten, a weist stabil Leeschtung no multiple Läsionen.
D'Dauer vun der Hash-Tabelle fir d'GPU an d'Interprogramm Kommunikatioun war 984 ms. Dëst beinhalt d'Zäit fir den Dësch an d'Erënnerung ze placéieren an ze läschen (allokéiert 1 GB Erënnerung eng Kéier, wat e bëssen Zäit an CUDA dauert), Elementer aginn an ze läschen, an iwwer hinnen iteréieren. All Kopie vun an aus der Video Kaart Erënnerung sinn och Rechnung gedroe.
Den Hash-Tabelle selwer huet 271 ms gedauert fir ze kompletéieren. Dëst beinhalt d'Zäit, déi d'Videokaart verbréngt fir Elementer ze setzen an ze läschen, a berücksichtegt net d'Zäit, déi d'Kopie an d'Erënnerung an d'Iteratioun iwwer déi resultéierend Tabell verbraucht. Wann d'GPU Dësch fir eng laang Zäit lieft, oder wann den Hash Dësch ganz an der Erënnerung vun der Video Kaart enthale ass (zum Beispill, fir en Hash Dësch ze kreéieren deen duerch aner GPU Code benotzt gëtt an net den zentrale Prozessor), dann d'Testresultat ass relevant.
Den Hash-Table fir eng Videokaart weist héich Leeschtung wéinst héijen Duerchgang an aktiver Paralleliséierung.
Defiziter
D'Hash Tabellarchitektur huet e puer Themen fir bewosst ze sinn:
- Linear Sonde gëtt duerch Clustering behënnert, wat verursaacht datt d'Schlësselen an der Tabell manner wéi perfekt plazéiert sinn.
- Schlësselen ginn net mat der Funktioun ewechgeholl
deletean iwwer Zäit clutter se den Dësch.
Als Resultat kann d'Performance vun engem Hash-Tabelle graduell degradéieren, besonnesch wann et fir eng laang Zäit existéiert an vill Inserts a Läschen huet. Ee Wee fir dës Nodeeler ze reduzéieren ass an en neien Dësch mat engem zimlech nidderegen Notzungsquote ze rehashen an déi ewechgeholl Schlësselen wärend der Rehashing auszefilteren.
Fir déi beschriwwe Themen ze illustréieren, benotzen ech den uewe genannte Code fir en Dësch mat 128 Milliounen Elementer ze kreéieren an duerch 4 Milliounen Elementer ze schloen bis ech 124 Millioune Plaze gefëllt hunn (Notzungsquote vu ronn 0,96). Hei ass d'Resultat Tabelle, all Zeil ass e CUDA Kernel Uruff fir 4 Milliounen nei Elementer an eng Hash Tabelle anzeginn:
Notzung Taux
Aféierungsdauer 4 Elementer
0,00
11,608448 ms (361,314798 Millioune Schlësselen/sec.)
0,03
11,751424 ms (356,918799 Millioune Schlësselen/sec.)
0,06
11,942592 ms (351,205515 Millioune Schlësselen/sec.)
0,09
12,081120 ms (347,178429 Millioune Schlësselen/sec.)
0,12
12,242560 ms (342,600233 Millioune Schlësselen/sec.)
0,16
12,396448 ms (338,347235 Millioune Schlësselen/sec.)
0,19
12,533024 ms (334,660176 Millioune Schlësselen/sec.)
0,22
12,703328 ms (330,173626 Millioune Schlësselen/sec.)
0,25
12,884512 ms (325,530693 Millioune Schlësselen/sec.)
0,28
13,033472 ms (321,810182 Millioune Schlësselen/sec.)
0,31
13,239296 ms (316,807174 Millioune Schlësselen/sec.)
0,34
13,392448 ms (313,184256 Millioune Schlësselen/sec.)
0,37
13,624000 ms (307,861434 Millioune Schlësselen/sec.)
0,41
13,875520 ms (302,280855 Millioune Schlësselen/sec.)
0,44
14,126528 ms (296,909756 Millioune Schlësselen/sec.)
0,47
14,399328 ms (291,284699 Millioune Schlësselen/sec.)
0,50
14,690304 ms (285,515123 Millioune Schlësselen/sec.)
0,53
15,039136 ms (278,892623 Millioune Schlësselen/sec.)
0,56
15,478656 ms (270,973402 Millioune Schlësselen/sec.)
0,59
15,985664 ms (262,379092 Millioune Schlësselen/sec.)
0,62
16,668673 ms (251,627968 Millioune Schlësselen/sec.)
0,66
17,587200 ms (238,486174 Millioune Schlësselen/sec.)
0,69
18,690048 ms (224,413765 Millioune Schlësselen/sec.)
0,72
20,278816 ms (206,831789 Millioune Schlësselen/sec.)
0,75
22,545408 ms (186,038058 Millioune Schlësselen/sec.)
0,78
26,053312 ms (160,989275 Millioune Schlësselen/sec.)
0,81
31,895008 ms (131,503463 Millioune Schlësselen/sec.)
0,84
42,103294 ms (99,619378 Millioune Schlësselen/sec.)
0,87
61,849056 ms (67,815164 Millioune Schlësselen/sec.)
0,90
105,695999 ms (39,682713 Millioune Schlësselen/sec.)
0,94
240,204636 ms (17,461378 Millioune Schlësselen/sec.)
Wéi d'Notzung eropgeet, geet d'Performance erof. Dëst ass an de meeschte Fäll net wënschenswäert. Wann eng Applikatioun Elementer an eng Tabell setzt an se dann verwerft (zum Beispill wann Dir Wierder an engem Buch zielt), dann ass dëst kee Problem. Awer wann d'Applikatioun eng laanglieweg Hash-Tabelle benotzt (zum Beispill an engem Grafikeditor fir net eidel Deeler vu Biller ze späicheren, wou de Benotzer dacks Informatioun setzt an läscht), da kann dëst Verhalen problematesch sinn.
A gemooss den Hash-Tabel-Sonddéift no 64 Milliounen Inserts (Notzungsfaktor 0,5). D'Duerchschnëttsdéift war 0,4774, sou datt déi meescht Schlësselen entweder am beschtméigleche Slot waren oder ee Slot ewech vun der beschter Positioun. Déi maximal Klangdéift war 60.
Ech hunn dunn d'Sonddéift op engem Dësch mat 124 Milliounen Inserts gemooss (Notzungsfaktor 0,97). Déi duerchschnëttlech Déift war schonn 10,1757, an déi maximal - 6474 (!!). Linear Sensing Leeschtung fällt wesentlech bei héijen Notzungsraten.
Et ass am beschten d'Notzungsquote vun dësem Hash-Tabelle niddereg ze halen. Awer dann erhéijen mir d'Performance op Käschte vum Erënnerungsverbrauch. Glécklecherweis, am Fall vun 32-Bit Schlësselen a Wäerter, kann dëst gerechtfäerdegt ginn. Wann am uewe genannte Beispill, an enger Tabell mat 128 Milliounen Elementer, mir den Notzungsfaktor vun 0,25 behalen, da kënne mir net méi wéi 32 Milliounen Elementer dra setzen, an déi reschtlech 96 Millioune Slots ginn verluer - 8 Bytes fir all Paar , 768 MB verluer Erënnerung.
Notéiert w.e.g. datt mir iwwer de Verloscht vu Videokartespeicher schwätzen, wat eng méi wäertvoll Ressource ass wéi Systemspeicher. Obwuel déi meescht modern Desktop Grafiken Kaarte datt CUDA Ënnerstëtzung hunn op d'mannst 4 GB vun Erënnerung (zu Schreiwen Zäit huet NVIDIA 2080 Ti 11 GB), et wier nach net déi schlau Decisioun esou Quantitéiten ze verléieren.
Méi spéit wäert ech méi iwwer d'Schafe vun Hash-Tabellen fir Videokaarten schreiwen, déi keng Probleemer mat der Ermëttlungsdéift hunn, souwéi Weeër fir geläscht Slots ze benotzen.
Kläng Déift Miessung
Fir d'Sondedéift vun engem Schlëssel ze bestëmmen, kënne mir den Hash vum Schlëssel (säin idealen Tabellindex) aus sengem aktuellen Tabellindex extrahéieren:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Wéinst der Magie vun zwee d'zwee komplementar binär Zuelen an der Tatsaach, datt d'Kapazitéit vun der Hash Dësch zwee zu der Muecht vun zwee ass, wäert dës Approche Aarbecht och wann de Schlëssel Index un den Ufank vun der Tabell geplënnert ass. Loosst eis e Schlëssel huelen, deen op 1 hashed ass, awer an de Slot gesat gëtt 3. Dann fir en Dësch mat Kapazitéit 4 kréien mir (3 — 1) & 3, wat entsprécht 2.
Konklusioun
Wann Dir Froen oder Kommentaren hutt, schéckt mech w.e.g. per E-Mail op oder en neit Thema opmaachen an.
Dëse Code gouf ënner Inspiratioun aus exzellenten Artikelen geschriwwen:
An Zukunft wäert ech weider iwwer Hash-Table Implementatiounen fir Videokaarten schreiwen an hir Leeschtung analyséieren. Meng Pläng enthalen chaining, Robin Hood hashing, a cuckoo hashing mat atomarer Operatiounen an Datestrukturen déi GPU frëndlech sinn.
Source: will.com