Feltettem a Github-ra .
Ez egy egyszerű GPU hash tábla, amely másodpercenként több száz millió beillesztést képes feldolgozni. Az NVIDIA GTX 1060 laptopomon a kód 64 millió véletlenszerűen generált kulcs-érték párt szúr be körülbelül 210 ms alatt, és eltávolít 32 millió párt körülbelül 64 ms alatt.
Ez azt jelenti, hogy egy laptop sebessége hozzávetőlegesen 300 millió beillesztés/sec és 500 millió törlés/sec.
A táblázat CUDA-ban van írva, bár ugyanez a technika alkalmazható HLSL-re vagy GLSL-re is. A megvalósításnak számos korlátozása van a videokártya nagy teljesítményének biztosítása érdekében:
- Csak a 32 bites kulcsok és ugyanazok az értékek kerülnek feldolgozásra.
- A hash tábla fix méretű.
- És ennek a méretnek egyenlőnek kell lennie a teljesítmény kettővel.
A kulcsokhoz és értékekhez le kell foglalni egy egyszerű határoló jelölőt (a fenti kódban ez 0xffffffff).
Hash asztal zárak nélkül
A hash tábla nyílt címzést használ a -val , azaz egyszerűen kulcs-érték párok tömbje, amely a memóriában van tárolva, és kiváló gyorsítótárteljesítménnyel rendelkezik. Ugyanez nem mondható el a láncolásról, amely magában foglalja a mutató keresését egy linkelt listában. A hash tábla egy egyszerű tömb, amely elemeket tárol KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
A táblázat mérete kettő hatványa, nem prímszám, mert a pow2/AND maszk alkalmazásához elég egy gyors utasítás, de a modulus operátor sokkal lassabb. Ez lineáris tapintás esetén fontos, mivel a lineáris táblázatkeresésben a slot indexet minden slotba be kell csomagolni. Ennek eredményeként a művelet költsége minden nyílásban modulo-val bővül.
A táblázat csak az egyes elemek kulcsát és értékét tárolja, a kulcs hash-jét nem. Mivel a táblázat csak 32 bites kulcsokat tárol, a hash kiszámítása nagyon gyorsan történik. A fenti kód a Murmur3 hash-t használja, amely csak néhány eltolást, XOR-t és szorzást hajt végre.
A hash tábla zárolásvédelmi technikákat használ, amelyek függetlenek a memória sorrendjétől. Még ha egyes írási műveletek megzavarják is a többi ilyen művelet sorrendjét, a hash tábla továbbra is fenntartja a megfelelő állapotot. Az alábbiakban erről fogunk beszélni. A technika kiválóan működik olyan videokártyákkal, amelyek több ezer szálat futnak egyszerre.
A hash táblában lévő kulcsok és értékek üresre vannak inicializálva.
A kód módosítható a 64 bites kulcsok és értékek kezelésére is. A kulcsokhoz atomi olvasási, írási és összehasonlítási és csereműveletre van szükség. Az értékekhez pedig atomi olvasási és írási műveletek szükségesek. Szerencsére a CUDA-ban a 32 és 64 bites értékek olvasási-írási műveletei atomi mindaddig, amíg természetesen igazodnak (lásd alább). ), a modern videokártyák pedig támogatják a 64 bites atomi összehasonlítási és csereműveleteket. Természetesen, ha 64 bitre váltunk, a teljesítmény kissé csökken.
Kivonattábla állapota
A hash táblában minden kulcs-érték párnak négy állapota lehet:
- A kulcs és az érték üres. Ebben az állapotban a hash tábla inicializálva van.
- A kulcsot felírták, de az értéket még nem írták le. Ha egy másik szál éppen olvas adatokat, akkor üresen tér vissza. Ez normális, ugyanez történt volna, ha egy másik végrehajtási szál kicsit korábban működött volna, és egy párhuzamos adatszerkezetről beszélünk.
- A kulcs és az érték is rögzítésre kerül.
- Az érték elérhető más végrehajtási szálak számára, de a kulcs még nem. Ez azért fordulhat elő, mert a CUDA programozási modellnek lazán rendezett memóriamodellje van. Ez normális; mindenesetre a kulcs továbbra is üres, még akkor is, ha az érték már nem az.
Fontos árnyalat, hogy miután a kulcsot beírtuk a slotba, az már nem mozdul – hiába törlik a kulcsot, erről lentebb lesz szó.
A hash tábla kódja még lazán rendezett memóriamodellekkel is működik, amelyekben a memória olvasási és írási sorrendje ismeretlen. Miközben a beszúrást, a keresést és a törlést nézzük egy hash-táblázatban, ne feledje, hogy minden kulcs-érték pár a fent leírt négy állapot valamelyikében van.
Beszúrás egy hash táblába
A CUDA függvény, amely kulcs-érték párokat szúr be egy hash táblába, így néz ki:
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);
}
}
Kulcs beszúrásához a kód a hash tábla tömbjén keresztül iterál, a beillesztett kulcs hashével kezdve. A tömb minden egyes helye atomi összehasonlítási és csereműveletet hajt végre, amely összehasonlítja az adott résben lévő kulcsot az üreshez. Ha a rendszer eltérést észlel, a nyílásban lévő kulcsot frissíti a behelyezett kulccsal, majd visszaküldi az eredeti kulcsot. Ha ez az eredeti kulcs üres volt, vagy megegyezett a behelyezett kulccsal, akkor a kód megfelelő helyet talált a beillesztéshez, és a beszúrt értéket beilleszti a nyílásba.
Ha egy kernelhívásban gpu_hashtable_insert() több elem van ugyanazzal a kulccsal, akkor ezek bármelyik értéke beírható a kulcsnyílásba. Ez normálisnak tekinthető: a hívás során valamelyik kulcsérték-írás sikerülni fog, de mivel mindez több végrehajtási szálon belül párhuzamosan történik, nem tudjuk megjósolni, hogy melyik memóriaírás lesz az utolsó.
Hash tábla keresése
Kód a kulcsok kereséséhez:
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);
}
}
Egy táblában tárolt kulcs értékének meghatározásához iteráljuk a tömböt a keresett kulcs hashével kezdve. Minden slotban ellenőrizzük, hogy a kulcs az, amit keresünk, és ha igen, akkor visszaadjuk az értékét. Azt is ellenőrizzük, hogy a kulcs üres-e, és ha igen, megszakítjuk a keresést.
Ha nem találjuk a kulcsot, a kód üres értéket ad vissza.
Mindezek a keresési műveletek egyidejűleg is végrehajthatók beszúrással és törléssel. A táblázatban minden pár rendelkezik a fent leírt négy állapot valamelyikével.
Törlés hash táblában
A kulcsok törlésének kódja:
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 kulcs törlése szokatlan módon történik: a kulcsot a táblázatban hagyjuk, és az értékét (nem magát a kulcsot) üresnek jelöljük. Ez a kód nagyon hasonlít a lookup(), kivéve, hogy ha egy kulcsot találunk, akkor az értéke üres lesz.
Ahogy fentebb említettük, ha egy kulcsot egy nyílásba írnak, az többé nem kerül áthelyezésre. Még ha egy elemet törölnek is a táblázatból, a kulcs a helyén marad, értéke egyszerűen üres lesz. Ez azt jelenti, hogy nem kell atomi írási műveletet használnunk a slot értékéhez, mert nem mindegy, hogy az aktuális érték üres-e vagy sem - attól még üres lesz.
Hash táblázat átméretezése
A hash tábla méretét úgy módosíthatja, hogy létrehoz egy nagyobb táblát, és a régi tábla nem üres elemeit illeszti be abba. Nem valósítottam meg ezt a funkciót, mert a mintakódot egyszerű akartam tartani. Ezenkívül a CUDA programokban a memóriafoglalás gyakran a gazdagép kódban történik, nem pedig a CUDA kernelben.
A cikk leírja, hogyan lehet módosítani egy ilyen zárral védett adatstruktúrát.
Versenyképesség
A fenti függvénykódrészletekben gpu_hashtable_insert(), _lookup() и _delete() egyszerre csak egy kulcs-érték párt dolgozzon fel. És lejjebb gpu_hashtable_insert(), _lookup() и _delete() párhuzamosan dolgozzon fel egy pár tömböt, mindegyik pár külön GPU-végrehajtási szálban:
// 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 zárolásálló hash-tábla támogatja az egyidejű beszúrásokat, kereséseket és törléseket. Mivel a kulcs-érték párok mindig a négy állapot valamelyikében vannak, és a kulcsok nem mozognak, a táblázat garantálja a helyességet, még akkor is, ha különböző típusú műveleteket használnak egyszerre.
Ha azonban párhuzamosan dolgozunk fel egy köteg beillesztést és törlést, és ha a párok bemeneti tömbje duplikált kulcsokat tartalmaz, akkor nem tudjuk megjósolni, hogy melyik pár „nyer” – az utolsóként kerül beírásra a hash táblába. Tegyük fel, hogy a beillesztési kódot párok bemeneti tömbjével hívtuk meg A/0 B/1 A/2 C/3 A/4. Amikor a kód befejeződött, párosítás B/1 и C/3 garantáltan jelen lesznek a táblázatban, de ugyanakkor bármelyik pár megjelenik benne A/0, A/2 vagy A/4. Ez lehet probléma, de lehet, hogy nem – minden az alkalmazástól függ. Előre tudhatja, hogy a bemeneti tömbben nincsenek duplikált kulcsok, vagy nem érdekli, melyik érték íródott utoljára.
Ha ez problémát okoz Önnek, akkor a duplikált párokat szét kell választania különböző CUDA-rendszerhívásokra. A CUDA-ban minden kernelt meghívó művelet mindig a következő kernelhívás előtt befejeződik (legalább egy szálon belül. A különböző szálakban a kernelek párhuzamosan futnak). A fenti példában, ha meghív egy kernelt a A/0 B/1 A/2 C/3, a másik pedig A/4, majd a kulcsot A megkapja az értéket 4.
Most beszéljünk arról, hogy a függvényeknek kell-e lookup() и delete() használjon sima vagy illékony mutatót a hash táblában lévő párok tömbjére. Azt állítja:
A fordító dönthet úgy, hogy optimalizálja az olvasást és írást a globális vagy megosztott memóriába... Ezek az optimalizálások a kulcsszó használatával letilthatók
volatile: ... minden erre a változóra való hivatkozás valódi memória olvasási vagy írási utasításba kerül lefordításra.
A helyességi megfontolások nem igényelnek alkalmazást volatile. Ha a végrehajtási szál egy korábbi olvasási művelet gyorsítótárazott értékeit használja, akkor kissé elavult információkat fog használni. De mégis, ez az információ a hash tábla helyes állapotából a kernelhívás egy bizonyos pillanatában. Ha a legfrissebb információkat kell használnia, használhatja az indexet volatile, de ekkor enyhén csökken a teljesítmény: tesztjeim szerint 32 millió elem törlésekor a sebesség 500 millió törlés/sec-ről 450 millió törlés/sec-re csökkent.
termelékenység
A 64 millió elem beillesztésére és 32 millió törlésére vonatkozó tesztben verseny alakult ki között std::unordered_map és gyakorlatilag nincs hash tábla a GPU-hoz:
std::unordered_map 70 691 ms-ot töltött az elemek behelyezésével és eltávolításával, majd kiszabadításával unordered_map (több millió elemtől való megszabadulás sok időt vesz igénybe, mert belül unordered_map több memóriafoglalás történik). Őszintén szólva, std:unordered_map teljesen más korlátozások. Ez egy egyetlen CPU-végrehajtási szál, bármilyen méretű kulcsértéket támogat, jól teljesít magas kihasználtsági arány mellett, és több törlés után is stabil teljesítményt mutat.
A GPU és a programok közötti kommunikáció hash táblájának időtartama 984 ms volt. Ez magában foglalja a táblázat memóriába helyezésével és törlésével (egyszeri 1 GB memória lefoglalásával, ami CUDA-ban némi időt vesz igénybe), az elemek beszúrásával és törlésével, valamint a felettük való iterációval töltött időt. A videokártya memóriájába érkező és onnan származó összes másolat is figyelembe lesz véve.
Maga a hash tábla 271 ms-ig tartott. Ez magában foglalja a videokártya által az elemek beillesztésével és törlésével töltött időt, és nem veszi figyelembe a memóriába másolás és a kapott táblázatban való iteráció idejét. Ha a GPU tábla hosszú ideig él, vagy ha a hash tábla teljes egészében a videokártya memóriájában van (például egy olyan hash tábla létrehozásához, amelyet más GPU kód, és nem a központi processzor használ), akkor a teszt eredménye releváns.
A videokártya hash táblázata nagy teljesítményt mutat a nagy áteresztőképesség és az aktív párhuzamosítás miatt.
Korlátozások
A hash tábla architektúrának van néhány problémája, amelyekkel tisztában kell lenni:
- A lineáris szondázást nehezíti a klaszterezés, ami azt eredményezi, hogy a billentyűk a táblázatban nem megfelelően helyezkednek el.
- A gombok nem távolíthatók el a funkció használatával
deleteés idővel összeborítják az asztalt.
Ennek eredményeként a hash tábla teljesítménye fokozatosan romolhat, különösen, ha hosszú ideig létezik, és számos beszúrást és törlést tartalmaz. E hátrányok mérséklésének egyik módja az, hogy egy új, meglehetősen alacsony kihasználtságú táblázatot készítünk, és az újrafeldolgozás során kiszűrjük az eltávolított kulcsokat.
A leírt problémák szemléltetésére a fenti kóddal létrehozok egy 128 millió elemet tartalmazó táblázatot, és végigfutok 4 millió elemen, amíg ki nem töltöm a 124 millió helyet (a kihasználtság körülbelül 0,96). Íme az eredménytábla, minden sor egy CUDA kernelhívás, amely 4 millió új elemet illeszt be egy hash-táblázatba:
Felhasználási arány
A beillesztés időtartama 4 194 304 elem
0,00
11,608448 ms (361,314798 millió kulcs/mp)
0,03
11,751424 ms (356,918799 millió kulcs/mp)
0,06
11,942592 ms (351,205515 millió kulcs/mp)
0,09
12,081120 ms (347,178429 millió kulcs/mp)
0,12
12,242560 ms (342,600233 millió kulcs/mp)
0,16
12,396448 ms (338,347235 millió kulcs/mp)
0,19
12,533024 ms (334,660176 millió kulcs/mp)
0,22
12,703328 ms (330,173626 millió kulcs/mp)
0,25
12,884512 ms (325,530693 millió kulcs/mp)
0,28
13,033472 ms (321,810182 millió kulcs/mp)
0,31
13,239296 ms (316,807174 millió kulcs/mp)
0,34
13,392448 ms (313,184256 millió kulcs/mp)
0,37
13,624000 ms (307,861434 millió kulcs/mp)
0,41
13,875520 ms (302,280855 millió kulcs/mp)
0,44
14,126528 ms (296,909756 millió kulcs/mp)
0,47
14,399328 ms (291,284699 millió kulcs/mp)
0,50
14,690304 ms (285,515123 millió kulcs/mp)
0,53
15,039136 ms (278,892623 millió kulcs/mp)
0,56
15,478656 ms (270,973402 millió kulcs/mp)
0,59
15,985664 ms (262,379092 millió kulcs/mp)
0,62
16,668673 ms (251,627968 millió kulcs/mp)
0,66
17,587200 ms (238,486174 millió kulcs/mp)
0,69
18,690048 ms (224,413765 millió kulcs/mp)
0,72
20,278816 ms (206,831789 millió kulcs/mp)
0,75
22,545408 ms (186,038058 millió kulcs/mp)
0,78
26,053312 ms (160,989275 millió kulcs/mp)
0,81
31,895008 ms (131,503463 millió kulcs/mp)
0,84
42,103294 ms (99,619378 millió kulcs/mp)
0,87
61,849056 ms (67,815164 millió kulcs/mp)
0,90
105,695999 ms (39,682713 millió kulcs/mp)
0,94
240,204636 ms (17,461378 millió kulcs/mp)
A kihasználtság növekedésével a teljesítmény csökken. Ez a legtöbb esetben nem kívánatos. Ha egy alkalmazás elemeket szúr be egy táblázatba, majd elveti azokat (például egy könyvben lévő szavak számlálásakor), akkor ez nem probléma. De ha az alkalmazás hosszú élettartamú hash-táblázatot használ (például egy grafikus szerkesztőben a képek nem üres részeit tárolja, ahol a felhasználó gyakran szúr be és töröl információkat), akkor ez a viselkedés problémás lehet.
És megmérte a hash táblázat szondázási mélységét 64 millió betét után (kihasználási tényező 0,5). Az átlagos mélység 0,4774 volt, tehát a legtöbb billentyű vagy a lehető legjobb nyílásban volt, vagy egy nyílással távolabb volt a legjobb pozíciótól. A maximális szondázási mélység 60 volt.
Ezután megmértem a tapintási mélységet egy asztalon, 124 millió lapkával (kihasználtsági tényező 0,97). Az átlagos mélység már 10,1757 volt, a maximális pedig - 6474 (!!). A lineáris érzékelési teljesítmény jelentősen csökken magas kihasználtság mellett.
A legjobb, ha ennek a hash-táblának a kihasználtságát alacsonyan tartja. Ekkor azonban a memóriafelhasználás rovására növeljük a teljesítményt. Szerencsére a 32 bites kulcsok és értékek esetében ez indokolható. Ha a fenti példában egy 128 millió elemet tartalmazó táblázatban a 0,25-ös kihasználtsági tényezőt megtartjuk, akkor legfeljebb 32 millió elemet helyezhetünk el benne, és a maradék 96 millió hely elveszik - páronként 8 bájt , 768 MB elveszett memória.
Felhívjuk figyelmét, hogy a videokártya memória elvesztéséről beszélünk, amely értékesebb erőforrás, mint a rendszermemória. Bár a legtöbb modern asztali grafikus kártya, amely támogatja a CUDA-t, legalább 4 GB memóriával rendelkezik (a cikk írásakor az NVIDIA 2080 Ti 11 GB-os), mégsem lenne a legbölcsebb döntés ekkora mennyiséget elveszíteni.
Később bővebben írok a hash-táblázatok létrehozásáról olyan videokártyákhoz, amelyeknek nincs problémája a szondázási mélységgel, valamint a törölt helyek újrafelhasználásának módjairól.
Szondázási mélységmérés
Egy kulcs vizsgálati mélységének meghatározásához kivonhatjuk a kulcs hash-jét (ideális táblázatindexét) a tényleges táblaindexből:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
A kettő kettő komplementer bináris számainak varázsereje és az a tény, hogy a hash tábla kapacitása kettő hatványa, ez a megközelítés akkor is működik, ha a kulcsindexet a táblázat elejére helyezzük. Vegyünk egy kulcsot, amely 1-esre van kivonatolva, de be van szúrva a 3-as nyílásba. Ekkor egy 4-es kapacitású táblára azt kapjuk, hogy (3 — 1) & 3, ami egyenértékű a 2-vel.
Következtetés
Ha kérdése vagy észrevétele van, kérjük, írjon nekem a következő e-mail címre vagy nyisson új témát.
Ez a kód kiváló cikkek ihletésére készült:
A jövőben továbbra is fogok írni a videokártyák hash tábla implementációiról, és elemzem a teljesítményüket. Terveim között szerepel a láncolás, a Robin Hood-kivonat és a kakukk-kivonat atomi műveletek segítségével olyan adatstruktúrákban, amelyek GPU-barátak.
Forrás: will.com