Ek het dit op Github geplaas .
Dit is 'n eenvoudige GPU-hash-tabel wat honderde miljoene insetsels per sekonde kan verwerk. Op my NVIDIA GTX 1060-skootrekenaar voeg die kode 64 miljoen willekeurig gegenereerde sleutel-waarde-pare in ongeveer 210 ms in en verwyder 32 miljoen pare in ongeveer 64 ms.
Dit wil sê, die spoed op 'n skootrekenaar is ongeveer 300 miljoen invoegings/sek en 500 miljoen deletes/sek.
Die tabel is in CUDA geskryf, alhoewel dieselfde tegniek op HLSL of GLSL toegepas kan word. Die implementering het verskeie beperkings om hoë werkverrigting op 'n videokaart te verseker:
- Slegs 32-bis sleutels en dieselfde waardes word verwerk.
- Die hash-tafel het 'n vaste grootte.
- En hierdie grootte moet gelyk wees aan twee tot die mag.
Vir sleutels en waardes moet jy 'n eenvoudige skeidingsmerker reserveer (in die bogenoemde kode is dit 0xffffffff).
Hash tafel sonder slotte
Die hash-tabel gebruik oop adressering met , dit wil sê, dit is bloot 'n reeks sleutel-waarde-pare wat in die geheue gestoor word en het voortreflike kasprestasie. Dieselfde kan nie gesê word vir ketting nie, wat die soeke na 'n wyser in 'n gekoppelde lys behels. 'n Hash-tabel is 'n eenvoudige skikking wat elemente stoor KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Die grootte van die tabel is 'n mag van twee, nie 'n priemgetal nie, want een vinnige instruksie is genoeg om die pow2/AND-masker toe te pas, maar die modulusoperateur is baie stadiger. Dit is belangrik in die geval van lineêre ondersoek, aangesien in 'n lineêre tabel-opsoek die gleufindeks in elke gleuf toegedraai moet word. En as gevolg hiervan word die koste van die operasie modulo in elke gleuf bygevoeg.
Die tabel stoor slegs die sleutel en waarde vir elke element, nie 'n hash van die sleutel nie. Aangesien die tabel slegs 32-bis-sleutels stoor, word die hash baie vinnig bereken. Die kode hierbo gebruik die Murmur3-hash, wat slegs 'n paar skofte, XOR's en vermenigvuldiging uitvoer.
Die hash-tabel gebruik sluitbeskermingstegnieke wat onafhanklik is van geheuevolgorde. Selfs as sommige skryfbewerkings die volgorde van ander sulke bewerkings ontwrig, sal die hash-tabel steeds die korrekte toestand handhaaf. Ons sal hieronder hieroor praat. Die tegniek werk uitstekend met videokaarte wat duisende drade gelyktydig loop.
Die sleutels en waardes in die hash-tabel word geïnisialiseer om leeg te wees.
Die kode kan gewysig word om ook 64-bis sleutels en waardes te hanteer. Sleutels vereis atoomlees, skryf en vergelyk-en-ruil bewerkings. En waardes vereis atoomlees- en skryfbewerkings. Gelukkig, in CUDA, is lees-skryf-bewerkings vir 32- en 64-bis-waardes atoom solank dit natuurlik in lyn is (sien hieronder). ), en moderne videokaarte ondersteun 64-bis atoomvergelyk-en-uitruil-operasies. Natuurlik, wanneer na 64 bisse beweeg, sal werkverrigting effens afneem.
Hash tabel toestand
Elke sleutel-waarde-paar in 'n hash-tabel kan een van vier toestande hê:
- Sleutel en waarde is leeg. In hierdie toestand word die hash-tabel geïnitialiseer.
- Die sleutel is afgeskryf, maar die waarde is nog nie afgeskryf nie. As 'n ander draad tans data lees, keer dit leeg terug. Dit is normaal, dieselfde ding sou gebeur het as 'n ander draad van uitvoering 'n bietjie vroeër gewerk het, en ons praat van 'n gelyktydige datastruktuur.
- Beide die sleutel en die waarde word aangeteken.
- Die waarde is beskikbaar vir ander drade van uitvoering, maar die sleutel is nog nie. Dit kan gebeur omdat die CUDA-programmeringsmodel 'n losgeordende geheuemodel het. Dit is normaal; in elk geval is die sleutel nog leeg, selfs al is die waarde nie meer so nie.
'n Belangrike nuanse is dat sodra die sleutel na die gleuf geskryf is, dit nie meer beweeg nie - selfs al word die sleutel uitgevee, sal ons hieronder hieroor praat.
Die hash-tabelkode werk selfs met losgeordende geheuemodelle waarin die volgorde waarin geheue gelees en geskryf word, onbekend is. Terwyl ons kyk na invoeging, opsoek en verwydering in 'n hash-tabel, onthou dat elke sleutel-waarde-paar in een van die vier toestande hierbo beskryf is.
Invoeging in 'n hash-tabel
Die CUDA-funksie wat sleutel-waarde-pare in 'n hash-tabel invoeg, lyk soos volg:
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);
}
}
Om 'n sleutel in te voeg, herhaal die kode deur die hash-tabelskikking wat begin met die hash van die ingevoegde sleutel. Elke gleuf in die skikking voer 'n atoomvergelyk-en-ruil bewerking uit wat die sleutel in daardie gleuf met leeg vergelyk. As 'n wanpassing bespeur word, word die sleutel in die gleuf opgedateer met die ingevoegde sleutel, en dan word die oorspronklike gleufsleutel teruggestuur. As hierdie oorspronklike sleutel leeg was of ooreenstem met die ingevoegde sleutel, dan het die kode 'n geskikte gleuf vir invoeging gevind en die ingevoegde waarde in die gleuf ingevoeg.
As in een kernoproep gpu_hashtable_insert() daar is veelvuldige elemente met dieselfde sleutel, dan kan enige van hul waardes na die sleutelgleuf geskryf word. Dit word as normaal beskou: een van die sleutelwaarde-skryfwerk tydens die oproep sal slaag, maar aangesien dit alles parallel binne verskeie drade van uitvoering gebeur, kan ons nie voorspel watter geheueskryf die laaste een sal wees nie.
Hash-tabelopsoek
Kode vir soeksleutels:
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);
}
}
Om die waarde van 'n sleutel wat in 'n tabel gestoor is, te vind, herhaal ons deur die skikking wat begin met die hash van die sleutel waarna ons soek. In elke gleuf kyk ons of die sleutel die een is waarna ons soek, en indien wel, gee ons die waarde daarvan terug. Ons kyk ook of die sleutel leeg is, en indien wel, staak ons die soektog.
As ons nie die sleutel kan vind nie, gee die kode 'n leë waarde terug.
Al hierdie soekbewerkings kan gelyktydig uitgevoer word deur invoegings en skrapings. Elke paar in die tabel sal een van die vier toestande hê wat hierbo beskryf is vir die vloei.
Vee tans in 'n hash-tabel uit
Kode vir die uitvee van sleutels:
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);
}
}
Die uitvee van 'n sleutel word op 'n ongewone manier gedoen: ons laat die sleutel in die tabel en merk sy waarde (nie die sleutel self nie) as leeg. Hierdie kode is baie soortgelyk aan lookup(), behalwe dat wanneer 'n pasmaat op 'n sleutel gevind word, dit sy waarde leeg maak.
Soos hierbo genoem, sodra 'n sleutel na 'n gleuf geskryf is, word dit nie meer geskuif nie. Selfs wanneer 'n element uit die tabel geskrap word, bly die sleutel in plek, die waarde daarvan word eenvoudig leeg. Dit beteken dat ons nie 'n atoomskryfbewerking vir die gleufwaarde hoef te gebruik nie, want dit maak nie saak of die huidige waarde leeg is of nie - dit sal steeds leeg word.
Verander die grootte van 'n hash-tabel
Jy kan die grootte van 'n hash-tabel verander deur 'n groter tabel te skep en nie-leë elemente van die ou tabel daarin in te voeg. Ek het nie hierdie funksionaliteit geïmplementeer nie, want ek wou die voorbeeldkode eenvoudig hou. Boonop word geheuetoewysing in CUDA-programme dikwels in die gasheerkode gedoen eerder as in die CUDA-kern.
Die artikel beskryf hoe om so 'n slotbeskermde datastruktuur te wysig.
Mededingendheid
In die bogenoemde funksie kode brokkies gpu_hashtable_insert(), _lookup() и _delete() verwerk een sleutel-waarde-paar op 'n slag. En laer gpu_hashtable_insert(), _lookup() и _delete() verwerk 'n verskeidenheid pare parallel, elke paar in 'n aparte GPU-uitvoeringsdraad:
// 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);
}
}
Die slotbestande hash-tabel ondersteun gelyktydige invoegings, opsoeke en skrappings. Omdat sleutel-waarde-pare altyd in een van vier toestande is en die sleutels nie beweeg nie, waarborg die tabel korrektheid selfs wanneer verskillende tipes bewerkings gelyktydig gebruik word.
As ons egter 'n bondel invoegings en skrappings parallel verwerk, en as die invoerskikking van pare duplikaatsleutels bevat, sal ons nie kan voorspel watter pare sal "wen" nie—sal laaste na die hash-tabel geskryf word. Kom ons sê ons het die invoegingskode genoem met 'n invoerskikking van pare A/0 B/1 A/2 C/3 A/4. Wanneer die kode voltooi is, pare B/1 и C/3 is gewaarborg om in die tabel teenwoordig te wees, maar terselfdertyd sal enige van die pare daarin verskyn A/0, A/2 of A/4. Dit kan 'n probleem wees of nie - dit hang alles af van die toepassing. Jy weet dalk vooraf dat daar geen duplikaatsleutels in die invoerskikking is nie, of jy gee dalk nie om watter waarde laaste geskryf is nie.
As dit vir jou 'n probleem is, moet jy die duplikaatpare in verskillende CUDA-stelseloproepe skei. In CUDA voltooi enige bewerking wat die kern oproep, altyd voor die volgende kernoproep (ten minste binne een draad. In verskillende drade word pitte parallel uitgevoer). In die voorbeeld hierbo, as jy een kern noem met A/0 B/1 A/2 C/3, en die ander met A/4, dan die sleutel A die waarde sal kry 4.
Kom ons praat nou oor of funksies moet lookup() и delete() gebruik 'n gewone of vlugtige wyser na 'n verskeidenheid pare in die hash-tabel. Stel dat:
Die samesteller kan kies om lees en skryf na globale of gedeelde geheue te optimaliseer... Hierdie optimaliserings kan gedeaktiveer word deur die sleutelwoord te gebruik
volatile: ... enige verwysing na hierdie veranderlike word saamgestel in 'n werklike geheue lees of skryf instruksie.
Korrektheidsoorwegings vereis nie toepassing nie volatile. As die uitvoerdraad 'n kaswaarde van 'n vroeëre leesbewerking gebruik, sal dit effens verouderde inligting gebruik. Maar steeds, dit is inligting van die korrekte toestand van die hash-tabel op 'n sekere oomblik van die kernoproep. As jy die nuutste inligting moet gebruik, kan jy die indeks gebruik volatile, maar dan sal die werkverrigting effens afneem: volgens my toetse, wanneer 32 miljoen elemente uitgevee word, het die spoed van 500 miljoen skrappings/sek tot 450 miljoen skrappings/sek.
produktiwiteit
In die toets vir die invoeging van 64 miljoen elemente en die verwydering van 32 miljoen van hulle, kompetisie tussen std::unordered_map en daar is feitlik geen hash-tabel vir die GPU nie:
std::unordered_map het 70 691 ms spandeer om elemente in te voeg en te verwyder en dit dan te bevry unordered_map (Om van miljoene elemente ontslae te raak, neem baie tyd, want binne unordered_map veelvuldige geheuetoewysings gemaak word). Eerlik gesproke, std:unordered_map heeltemal ander beperkings. Dit is 'n enkele SVE-draad van uitvoering, ondersteun sleutelwaardes van enige grootte, presteer goed teen hoë benuttingskoerse en toon stabiele werkverrigting na veelvuldige skrappings.
Die tydsduur van die hash-tabel vir die GPU en inter-program kommunikasie was 984 ms. Dit sluit in die tyd wat spandeer word om die tabel in die geheue te plaas en dit uit te vee (toewysing van 1 GB geheue een keer, wat 'n geruime tyd in CUDA neem), die invoeging en uitvee van elemente, en die herhaling daarvan. Alle kopieë na en van die videokaartgeheue word ook in ag geneem.
Die hash-tabel self het 271 ms geneem om te voltooi. Dit sluit die tyd in wat die videokaart spandeer om elemente in te voeg en uit te vee, en neem nie die tyd in ag wat spandeer word om in die geheue te kopieer en oor die resulterende tabel te herhaal nie. As die GPU-tabel vir 'n lang tyd leef, of as die hash-tabel geheel en al in die geheue van die videokaart vervat is (byvoorbeeld om 'n hash-tabel te skep wat deur ander GPU-kode gebruik sal word en nie die sentrale verwerker nie), dan die toetsuitslag is relevant.
Die hash-tabel vir 'n videokaart toon hoë werkverrigting as gevolg van hoë deurset en aktiewe parallelisering.
Beperkings
Die hash-tabel-argitektuur het 'n paar probleme om van bewus te wees:
- Lineêre ondersoek word belemmer deur groepering, wat veroorsaak dat die sleutels in die tabel minder as perfek geplaas word.
- Sleutels word nie met die funksie verwyder nie
deleteen met verloop van tyd maak hulle die tafel deurmekaar.
Gevolglik kan die werkverrigting van 'n hash-tabel geleidelik verswak, veral as dit vir 'n lang tyd bestaan en talle invoegings en verwyderings het. Een manier om hierdie nadele te versag, is om na 'n nuwe tabel met 'n redelik lae benuttingskoers te herhash en die verwyderde sleutels uit te filter tydens die herhashing.
Om die kwessies wat beskryf word, te illustreer, sal ek die bogenoemde kode gebruik om 'n tabel met 128 miljoen elemente te skep en deur 4 miljoen elemente te loop totdat ek 124 miljoen gleuwe gevul het (benuttingskoers van ongeveer 0,96). Hier is die resultaattabel, elke ry is 'n CUDA-kernoproep om 4 miljoen nuwe elemente in een hash-tabel in te voeg:
Gebruikskoers
Invoegingsduur 4 194 304 elemente
0,00
11,608448 ms (361,314798 miljoen sleutels/sek.)
0,03
11,751424 ms (356,918799 miljoen sleutels/sek.)
0,06
11,942592 ms (351,205515 miljoen sleutels/sek.)
0,09
12,081120 ms (347,178429 miljoen sleutels/sek.)
0,12
12,242560 ms (342,600233 miljoen sleutels/sek.)
0,16
12,396448 ms (338,347235 miljoen sleutels/sek.)
0,19
12,533024 ms (334,660176 miljoen sleutels/sek.)
0,22
12,703328 ms (330,173626 miljoen sleutels/sek.)
0,25
12,884512 ms (325,530693 miljoen sleutels/sek.)
0,28
13,033472 ms (321,810182 miljoen sleutels/sek.)
0,31
13,239296 ms (316,807174 miljoen sleutels/sek.)
0,34
13,392448 ms (313,184256 miljoen sleutels/sek.)
0,37
13,624000 ms (307,861434 miljoen sleutels/sek.)
0,41
13,875520 ms (302,280855 miljoen sleutels/sek.)
0,44
14,126528 ms (296,909756 miljoen sleutels/sek.)
0,47
14,399328 ms (291,284699 miljoen sleutels/sek.)
0,50
14,690304 ms (285,515123 miljoen sleutels/sek.)
0,53
15,039136 ms (278,892623 miljoen sleutels/sek.)
0,56
15,478656 ms (270,973402 miljoen sleutels/sek.)
0,59
15,985664 ms (262,379092 miljoen sleutels/sek.)
0,62
16,668673 ms (251,627968 miljoen sleutels/sek.)
0,66
17,587200 ms (238,486174 miljoen sleutels/sek.)
0,69
18,690048 ms (224,413765 miljoen sleutels/sek.)
0,72
20,278816 ms (206,831789 miljoen sleutels/sek.)
0,75
22,545408 ms (186,038058 miljoen sleutels/sek.)
0,78
26,053312 ms (160,989275 miljoen sleutels/sek.)
0,81
31,895008 ms (131,503463 miljoen sleutels/sek.)
0,84
42,103294 ms (99,619378 miljoen sleutels/sek.)
0,87
61,849056 ms (67,815164 miljoen sleutels/sek.)
0,90
105,695999 ms (39,682713 miljoen sleutels/sek.)
0,94
240,204636 ms (17,461378 miljoen sleutels/sek.)
Soos benutting toeneem, neem prestasie af. Dit is in die meeste gevalle nie wenslik nie. As 'n toepassing elemente in 'n tabel invoeg en dit dan weggooi (byvoorbeeld wanneer woorde in 'n boek getel word), dan is dit nie 'n probleem nie. Maar as die toepassing 'n langlewende hash-tabel gebruik (byvoorbeeld in 'n grafiese redigeerder om nie-leë dele van beelde te stoor waar die gebruiker gereeld inligting invoeg en uitvee), dan kan hierdie gedrag problematies wees.
En gemeet die hash tabel peiling diepte na 64 miljoen insetsels (benutting faktor 0,5). Die gemiddelde diepte was 0,4774, so die meeste sleutels was óf in die beste moontlike gleuf óf een gleuf weg van die beste posisie. Die maksimum klankdiepte was 60.
Ek het toe die peildiepte gemeet op 'n tafel met 124 miljoen insetsels (benuttingsfaktor 0,97). Die gemiddelde diepte was reeds 10,1757, en die maksimum - 6474 (!!). Lineêre waarnemingsprestasie daal aansienlik teen hoë benuttingskoerse.
Dit is die beste om hierdie hash-tabel se gebruikskoers laag te hou. Maar dan verhoog ons werkverrigting ten koste van geheueverbruik. Gelukkig, in die geval van 32-bis sleutels en waardes, kan dit geregverdig word. As ons in die voorbeeld hierbo, in 'n tabel met 128 miljoen elemente, die benuttingsfaktor van 0,25 hou, dan kan ons nie meer as 32 miljoen elemente daarin plaas nie, en die oorblywende 96 miljoen gleuwe sal verlore gaan - 8 grepe vir elke paar , 768 MB verlore geheue.
Neem asseblief kennis dat ons praat oor die verlies van videokaartgeheue, wat 'n meer waardevolle hulpbron as stelselgeheue is. Alhoewel die meeste moderne rekenaargrafiese kaarte wat CUDA ondersteun, ten minste 4 GB geheue het (ten tyde van skryf het die NVIDIA 2080 Ti 11 GB), sou dit steeds nie die verstandigste besluit wees om sulke bedrae te verloor nie.
Later sal ek meer skryf oor die skep van hash-tabelle vir videokaarte wat nie probleme het met die ondersoek van diepte nie, asook maniere om geskrap gleuwe te hergebruik.
Klankdieptemeting
Om die ondersoekdiepte van 'n sleutel te bepaal, kan ons die sleutel se hash (sy ideale tabelindeks) uit sy werklike tabelindeks onttrek:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
As gevolg van die magie van twee se twee se komplementêre binêre getalle en die feit dat die kapasiteit van die hash-tabel twee tot die krag van twee is, sal hierdie benadering werk selfs wanneer die sleutelindeks na die begin van die tabel geskuif word. Kom ons neem 'n sleutel wat na 1 gehas het, maar in gleuf 3 geplaas word. Dan kry ons vir 'n tafel met kapasiteit 4 (3 — 1) & 3, wat gelykstaande is aan 2.
Gevolgtrekking
As jy vrae of opmerkings het, e-pos my asseblief by of maak 'n nuwe onderwerp oop in.
Hierdie kode is geskryf onder inspirasie van uitstekende artikels:
In die toekoms sal ek voortgaan om te skryf oor hash-tabelimplementerings vir videokaarte en hul werkverrigting te ontleed. My planne sluit in ketting, Robin Hood-hashing en koekoek-hashing met behulp van atoombewerkings in datastrukture wat GPU-vriendelik is.
Bron: will.com