Ég setti það á Github .
Þetta er einföld GPU kjötkássatafla sem getur unnið úr hundruð milljóna innskots á sekúndu. Á NVIDIA GTX 1060 fartölvunni minni setur kóðinn inn 64 milljón af handahófi mynduð lykilgildapör á um 210 ms og fjarlægir 32 milljónir pör á um 64 ms.
Það er að segja að hraðinn á fartölvu er um það bil 300 milljón innsetningar/sek og 500 milljónir eyðingar/sek.
Taflan er skrifuð í CUDA, þó hægt sé að beita sömu tækni á HLSL eða GLSL. Útfærslan hefur nokkrar takmarkanir til að tryggja mikla afköst á skjákorti:
- Aðeins 32 bita lyklar og sömu gildi eru unnin.
- Hash borðið er með fastri stærð.
- Og þessi stærð verður að vera jöfn tveimur í krafti.
Fyrir lykla og gildi þarftu að panta einfalt afmörkunarmerki (í ofangreindum kóða er þetta 0xffffffff).
Hash borð án lása
Hash taflan notar opna heimilisfang með , það er, það er einfaldlega fylki lykilgilda pöra sem er geymt í minni og hefur yfirburða skyndiminni. Það sama er ekki hægt að segja um chaining, sem felur í sér að leita að bendili á tengdum lista. Kjötkássatafla er einfalt fylki sem geymir þætti KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Stærð töflunnar er kraftur tveggja, ekki prímtala, því ein hröð leiðbeining nægir til að nota pow2/AND grímuna, en stuðullinn er mun hægari. Þetta er mikilvægt þegar um línulega könnun er að ræða, þar sem í línulegri töfluuppflettingu verður rifavísitalan að vera í hverri rauf. Og þar af leiðandi er kostnaður við aðgerðina bætt við modulo í hverri rauf.
Taflan geymir aðeins lykilinn og gildi fyrir hvern þátt, ekki kjötkássa af lyklinum. Þar sem taflan geymir aðeins 32 bita lykla reiknast kjötkássa mjög fljótt. Kóðinn hér að ofan notar Murmur3 kjötkássa, sem framkvæmir aðeins nokkrar breytingar, XOR og margföldun.
Hash-taflan notar læsingarverndaraðferðir sem eru óháðar minnisröð. Jafnvel þótt sumar skrifaðgerðir trufli röð annarra slíkra aðgerða mun kjötkássataflan samt halda réttu ástandi. Við munum tala um þetta hér að neðan. Tæknin virkar frábærlega með skjákortum sem keyra þúsundir þráða samtímis.
Lyklar og gildi í kjötkássatöflunni eru frumstillt til að vera tóm.
Hægt er að breyta kóðanum til að meðhöndla 64-bita lykla og gildi líka. Lyklar krefjast atómalesturs, ritunar og samanburðar-og-skipta. Og gildi krefjast atómlestrar- og ritunaraðgerða. Sem betur fer, í CUDA, eru les- og skrifaðgerðir fyrir 32- og 64-bita gildi frumeinda svo lengi sem þær eru náttúrulega samræmdar (sjá hér að neðan). ), og nútíma skjákort styðja 64-bita atóm samanburðar- og skiptiaðgerðir. Auðvitað, þegar þú færir yfir í 64 bita, mun frammistaða minnka lítillega.
Hash töflu ástand
Hvert lykilgilda par í kjötkássatöflu getur haft eitt af fjórum stöðum:
- Lykill og gildi eru tóm. Í þessu ástandi er kjötkássataflan frumstillt.
- Lykillinn hefur verið skráður niður, en gildið hefur ekki enn verið skrifað. Ef annar þráður er núna að lesa gögn, þá skilar hann tómum. Þetta er eðlilegt, það sama hefði gerst ef annar útfærsluþráður hefði virkað aðeins fyrr og við erum að tala um samhliða gagnaskipulag.
- Bæði lykillinn og gildið eru skráð.
- Gildið er í boði fyrir aðra útfærsluþræði, en lykillinn er ekki enn. Þetta getur gerst vegna þess að CUDA forritunarlíkanið er með lauslega raðað minnislíkan. Þetta er eðlilegt; í öllum tilvikum er lykillinn enn tómur, jafnvel þótt gildið sé það ekki lengur.
Mikilvægur blæbrigði er að þegar lykillinn hefur verið skrifaður í raufina hreyfist hann ekki lengur - jafnvel þótt lyklinum sé eytt munum við tala um þetta hér að neðan.
Kóðinn fyrir kjötkássatöflu virkar meira að segja með lauslega röðum minnislíkönum þar sem röð minnis er lesið og skrifað er óþekkt. Þegar við skoðum innsetningu, uppflettingu og eyðingu í kjötkássatöflu, mundu að hvert lykilgildi par er í einu af fjórum stöðum sem lýst er hér að ofan.
Að setja inn í kjötkássatöflu
CUDA aðgerðin sem setur lykilgildapör inn í kjötkássatöflu lítur svona út:
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);
}
}
Til að setja inn lykil fer kóðinn í gegnum kjötkássatöflufylkinguna og byrjar á kjötkássa lykilsins sem settur var inn. Hver rauf í fylkinu framkvæmir atómsamanburð-og-skiptaaðgerð sem ber saman lykilinn í þeirri rauf við tóman. Ef ósamræmi greinist er lykillinn í raufinni uppfærður með innsettum lykli og síðan er upprunalega raufarlyklinum skilað. Ef þessi upprunalegi lykill var tómur eða passaði við innsetta lykilinn, þá fann kóðinn viðeigandi rauf til innsetningar og setti innsetta gildið inn í raufina.
Ef í einu kjarnakalli gpu_hashtable_insert() það eru margir þættir með sama lyklinum, þá er hægt að skrifa hvaða gildi þeirra í lykilraufina. Þetta er talið eðlilegt: ein af lykilgildisritunum meðan á símtalinu stendur mun heppnast, en þar sem allt þetta gerist samhliða innan nokkurra þráða framkvæmdar, getum við ekki spáð fyrir um hvaða minnisritun verður sú síðasta.
Hash töfluleit
Kóði fyrir leitarlykla:
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);
}
}
Til að finna gildi lykils sem geymdur er í töflu, endurtekum við fylkið sem byrjar á kjötkássa lykilsins sem við erum að leita að. Í hverri rauf athugum við hvort lykillinn sé sá sem við erum að leita að og ef svo er þá skilum við gildi hans. Við athugum líka hvort lykillinn sé tómur og ef svo er hættum við leitinni.
Ef við finnum ekki lykilinn skilar kóðinn auðu gildi.
Allar þessar leitaraðgerðir er hægt að framkvæma samtímis með innsetningu og eyðingu. Hvert par í töflunni mun hafa eitt af fjórum ríkjum sem lýst er hér að ofan fyrir flæðið.
Eyðir í kjötkássatöflu
Kóði til að eyða lyklum:
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);
}
}
Að eyða lykli er gert á óvenjulegan hátt: við skiljum lykilinn eftir í töflunni og merkjum gildi hans (ekki lykilinn sjálfan) sem tóman. Þessi kóði er mjög svipaður lookup(), nema að þegar samsvörun finnst á lykli gerir það gildi hans tómt.
Eins og getið er hér að ofan, þegar lykill er skrifaður í rauf, er hann ekki lengur færður. Jafnvel þegar einingu er eytt úr töflunni er lykillinn áfram á sínum stað, gildi hans verður einfaldlega tómt. Þetta þýðir að við þurfum ekki að nota atómskrifaaðgerð fyrir rifagildið, því það skiptir ekki máli hvort núverandi gildi er tómt eða ekki - það verður samt tómt.
Breyta stærð kjötkássatöflu
Þú getur breytt stærð kjötkássatöflu með því að búa til stærri töflu og setja inn í hana ótóma þætti úr gömlu töflunni. Ég útfærði ekki þessa virkni vegna þess að ég vildi hafa sýnishornskóðann einfaldan. Þar að auki, í CUDA forritum, er minnisúthlutun oft gerð í hýsilkóðann frekar en í CUDA kjarnanum.
Í greininni lýsir því hvernig á að breyta slíkri læsingarvarinni gagnabyggingu.
Samkeppnishæfni
Í ofangreindum aðgerðakóðabútum gpu_hashtable_insert(), _lookup() и _delete() vinna úr einu lykilgildi pari í einu. Og lægri gpu_hashtable_insert(), _lookup() и _delete() vinna úr fjölda pöra samhliða, hvert par í sérstökum GPU-útfærsluþræði:
// 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);
}
}
Læsingarþolna kjötkássataflan styður samhliða innskot, uppflettingu og eyðingu. Þar sem lykilgildapör eru alltaf í einu af fjórum stöðum og lyklarnir hreyfast ekki, tryggir taflan réttmæti jafnvel þegar mismunandi gerðir aðgerða eru notaðar samtímis.
Hins vegar, ef við vinnum lotu af innsetningum og brottfellingum samhliða, og ef inntaksfylki pöranna inniheldur afrita lykla, þá getum við ekki spáð fyrir um hvaða pör munu „vinna“ - verður skrifað í kjötkássatöfluna síðast. Segjum að við kölluðum innsetningarkóðann með inntaksfylki af pörum A/0 B/1 A/2 C/3 A/4. Þegar kóðinn lýkur, parast B/1 и C/3 eru tryggð til staðar í töflunni, en á sama tíma mun eitthvað af pörunum birtast í henni A/0, A/2 eða A/4. Þetta gæti verið vandamál eða ekki - það veltur allt á forritinu. Þú gætir vitað fyrirfram að það eru engir tvíteknir lyklar í inntaksfylkingunni, eða þér er kannski sama hvaða gildi var skrifað síðast.
Ef þetta er vandamál fyrir þig, þá þarftu að aðgreina tvítekin pör í mismunandi CUDA kerfissímtöl. Í CUDA lýkur öllum aðgerðum sem kallar á kjarnann alltaf fyrir næsta kjarnakall (að minnsta kosti innan eins þráðs. Í mismunandi þráðum eru kjarna keyrðir samhliða). Í dæminu hér að ofan, ef þú kallar einn kjarna með A/0 B/1 A/2 C/3, og hinn með A/4, svo lykillinn A mun fá verðmæti 4.
Nú skulum við tala um hvort aðgerðir ættu lookup() и delete() notaðu látlausan eða óstöðugan bendil á fylki af pörum í kjötkássatöflunni. Segir að:
Þýðandinn gæti valið að fínstilla lestur og skrif í alþjóðlegt eða sameiginlegt minni... Þessar fínstillingar er hægt að gera óvirkar með því að nota lykilorðið
volatile: ... allar tilvísanir í þessa breytu eru settar saman í raunverulegt minni les- eða skrifleiðbeiningar.
Réttlætissjónarmið krefjast ekki umsóknar volatile. Ef framkvæmdarþráðurinn notar skyndiminni gildi frá fyrri lestri, þá mun hann nota aðeins úreltar upplýsingar. En samt eru þetta upplýsingar úr réttu ástandi kjötkássatöflunnar á ákveðnu augnabliki í kjarnakallinu. Ef þú þarft að nota nýjustu upplýsingarnar geturðu notað vísitöluna volatile, en þá mun frammistaðan minnka lítillega: samkvæmt prófunum mínum, þegar 32 milljón þáttum var eytt, lækkaði hraðinn úr 500 milljónum eyðingar/sek í 450 milljón eyðingar/sek.
Framleiðni
Í prófinu fyrir að setja inn 64 milljónir þátta og eyða 32 milljónum þeirra, samkeppni á milli std::unordered_map og það er nánast engin kjötkássatafla fyrir GPU:
std::unordered_map eyddi 70 ms í að setja inn og fjarlægja þætti og losa þá síðan unordered_map (að losa sig við milljónir frumefna tekur mikinn tíma, því að innan unordered_map margar minnisúthlutanir eru gerðar). Í hreinskilni sagt, std:unordered_map allt aðrar takmarkanir. Það er einn CPU-þráður framkvæmdar, styður lykilgildi af hvaða stærð sem er, skilar sér vel við háan nýtingarhraða og sýnir stöðugan árangur eftir margar eyðingar.
Lengd kjötkássatöflunnar fyrir GPU og samskipti milli forrita var 984 ms. Þetta felur í sér þann tíma sem fer í að setja töfluna í minni og eyða henni (úthluta 1 GB af minni einu sinni, sem tekur nokkurn tíma í CUDA), setja inn og eyða þáttum og endurtaka yfir þá. Öll afrit til og frá minni skjákortsins eru einnig tekin með í reikninginn.
Það tók 271 ms að klára kjötkássatöfluna sjálfa. Þetta felur í sér þann tíma sem skjákortið fer í að setja í og eyða þáttum, og tekur ekki tillit til þess tíma sem fer í að afrita í minni og endurtaka yfir töfluna sem myndast. Ef GPU borðið lifir í langan tíma, eða ef kjötkássataflan er að öllu leyti í minni skjákortsins (til dæmis til að búa til kjötkássatöflu sem verður notuð af öðrum GPU kóða en ekki miðlæga örgjörvanum), þá prófunarniðurstaðan skiptir máli.
Hash borðið fyrir skjákort sýnir mikla afköst vegna mikils afkösts og virkra samhliða.
Takmarkanir
Hash töfluarkitektúrinn hefur nokkur vandamál sem þarf að vera meðvitaður um:
- Línuleg könnun er hindruð af þyrping, sem veldur því að lyklarnir í töflunni eru ekki fullkomlega staðsettir.
- Lyklar eru ekki fjarlægðir með aðgerðinni
deleteog með tímanum ruglast þeir á borðinu.
Fyrir vikið getur frammistaða kjötkássatöflu minnkað smám saman, sérstaklega ef hún er til í langan tíma og hefur fjölmargar innsetningar og eyðingar. Ein leið til að draga úr þessum ókostum er að rehash inn í nýja töflu með frekar lágu nýtingarhlutfalli og sía út fjarlægðu lyklana meðan á endurhassuninni stendur.
Til að útskýra vandamálin sem lýst er, mun ég nota ofangreindan kóða til að búa til töflu með 128 milljón þáttum og fara í gegnum 4 milljónir þátta þar til ég hef fyllt 124 milljónir rifa (nýtingarhlutfall um 0,96). Hér er niðurstöðutaflan, hver röð er CUDA kjarnakall til að setja 4 milljónir nýrra þátta inn í eina kjötkássatöflu:
Notkunarhlutfall
Innsetningartími 4 þættir
0,00
11,608448 ms (361,314798 milljónir lykla/sek.)
0,03
11,751424 ms (356,918799 milljónir lykla/sek.)
0,06
11,942592 ms (351,205515 milljónir lykla/sek.)
0,09
12,081120 ms (347,178429 milljónir lykla/sek.)
0,12
12,242560 ms (342,600233 milljónir lykla/sek.)
0,16
12,396448 ms (338,347235 milljónir lykla/sek.)
0,19
12,533024 ms (334,660176 milljónir lykla/sek.)
0,22
12,703328 ms (330,173626 milljónir lykla/sek.)
0,25
12,884512 ms (325,530693 milljónir lykla/sek.)
0,28
13,033472 ms (321,810182 milljónir lykla/sek.)
0,31
13,239296 ms (316,807174 milljónir lykla/sek.)
0,34
13,392448 ms (313,184256 milljónir lykla/sek.)
0,37
13,624000 ms (307,861434 milljónir lykla/sek.)
0,41
13,875520 ms (302,280855 milljónir lykla/sek.)
0,44
14,126528 ms (296,909756 milljónir lykla/sek.)
0,47
14,399328 ms (291,284699 milljónir lykla/sek.)
0,50
14,690304 ms (285,515123 milljónir lykla/sek.)
0,53
15,039136 ms (278,892623 milljónir lykla/sek.)
0,56
15,478656 ms (270,973402 milljónir lykla/sek.)
0,59
15,985664 ms (262,379092 milljónir lykla/sek.)
0,62
16,668673 ms (251,627968 milljónir lykla/sek.)
0,66
17,587200 ms (238,486174 milljónir lykla/sek.)
0,69
18,690048 ms (224,413765 milljónir lykla/sek.)
0,72
20,278816 ms (206,831789 milljónir lykla/sek.)
0,75
22,545408 ms (186,038058 milljónir lykla/sek.)
0,78
26,053312 ms (160,989275 milljónir lykla/sek.)
0,81
31,895008 ms (131,503463 milljónir lykla/sek.)
0,84
42,103294 ms (99,619378 milljónir lykla/sek.)
0,87
61,849056 ms (67,815164 milljónir lykla/sek.)
0,90
105,695999 ms (39,682713 milljónir lykla/sek.)
0,94
240,204636 ms (17,461378 milljónir lykla/sek.)
Eftir því sem nýtingin eykst minnkar árangur. Þetta er ekki æskilegt í flestum tilfellum. Ef forrit setur þætti inn í töflu og fleygir þeim síðan (til dæmis þegar taldir eru orð í bók), þá er þetta ekki vandamál. En ef forritið notar langlífa kjötkássatöflu (til dæmis í grafíkritil til að geyma hluta mynda sem ekki eru tómir þar sem notandinn setur oft inn og eyðir upplýsingum), þá getur þessi hegðun verið erfið.
Og mældi kjötkássatöfluna til að rannsaka dýpt eftir 64 milljón innskot (nýtingarstuðull 0,5). Meðaldýptin var 0,4774, þannig að flestir lyklar voru annað hvort í bestu mögulegu raufinni eða einni rauf frá bestu stöðunni. Hámarks hljóðdýpt var 60.
Ég mældi svo könnunardýptina á töflu með 124 milljón innskotum (nýtingarstuðull 0,97). Meðaldýptin var þegar 10,1757 og hámarkið - 6474 (!!). Afköst línulegrar skynjunar lækkar verulega við háan nýtingarhraða.
Best er að hafa nýtingarhlutfall þessarar kjötkássatöflu lágt. En svo aukum við afköst á kostnað minnisnotkunar. Sem betur fer, ef um 32-bita lykla og gildi er að ræða, er hægt að réttlæta þetta. Ef í ofangreindu dæmi, í töflu með 128 milljón þáttum, höldum við nýtingarstuðlinum 0,25, þá getum við ekki sett meira en 32 milljónir þátta í það, og 96 milljónir rifa sem eftir eru tapast - 8 bæti fyrir hvert par , 768 MB glataðs minni.
Vinsamlegast athugaðu að við erum að tala um tap á skjákortaminni, sem er verðmætari auðlind en kerfisminni. Þó að flest nútíma skjákort sem styðja CUDA hafi að minnsta kosti 4 GB af minni (þegar þetta er skrifað hefur NVIDIA 2080 Ti 11 GB), þá væri það samt ekki skynsamlegasta ákvörðunin að tapa slíkum upphæðum.
Síðar mun ég skrifa meira um að búa til kjötkássatöflur fyrir skjákort sem eiga ekki í vandræðum með að rannsaka dýpt, sem og leiðir til að endurnýta eyddar raufar.
Hljómandi dýptarmæling
Til að ákvarða könnunardýpt lykils getum við dregið kjötkássa lykilsins (hugsjónatöfluvísitölu hans) úr raunverulegri töfluvísitölu hans:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Vegna töfra viðbótar tvöfaldra talna tveggja og þeirrar staðreyndar að getu kjötkássatöflunnar er tvö í krafti tveggja, mun þessi nálgun virka jafnvel þegar lykilvísitalan er færð í byrjun töflunnar. Tökum lykil sem hassaði í 1, en er settur í rauf 3. Síðan fáum við borð með rúmtak 4 (3 — 1) & 3, sem jafngildir 2.
Ályktun
Ef þú hefur spurningar eða athugasemdir, vinsamlegast sendu mér tölvupóst á eða opnaðu nýtt efni í.
Þessi kóða var skrifaður undir innblástur frá frábærum greinum:
Í framtíðinni mun ég halda áfram að skrifa um útfærslur á kjötkássatöflum fyrir skjákort og greina árangur þeirra. Áætlanir mínar fela í sér keðjutengingu, Robin Hood-kássingu og kúkakáss með því að nota atómaðgerðir í gagnamannvirkjum sem eru GPU-væn.
Heimild: www.habr.com