Zverejnil som to na Github .
Je to jednoduchá tabuľka hash GPU schopná spracovať stovky miliónov vložiek za sekundu. Na mojom notebooku NVIDIA GTX 1060 kód vloží 64 miliónov náhodne vygenerovaných párov kľúč-hodnota za približne 210 ms a odstráni 32 miliónov párov za približne 64 ms.
To znamená, že rýchlosť prenosného počítača je približne 300 miliónov vložení/s a 500 miliónov vymazaní/s.
Tabuľka je napísaná v CUDA, hoci rovnakú techniku možno použiť aj na HLSL alebo GLSL. Implementácia má niekoľko obmedzení na zabezpečenie vysokého výkonu na grafickej karte:
- Spracúvajú sa iba 32-bitové kľúče a rovnaké hodnoty.
- Tabuľka hash má pevnú veľkosť.
- A táto veľkosť sa musí rovnať dvom mocnine.
Pre kľúče a hodnoty si musíte rezervovať jednoduchý oddeľovač (vo vyššie uvedenom kóde je to 0xffffffff).
Hašovacia tabuľka bez zámkov
Tabuľka hash používa otvorené adresovanie s , to znamená, že je to jednoducho pole párov kľúč – hodnota, ktoré je uložené v pamäti a má vynikajúci výkon vyrovnávacej pamäte. To isté sa nedá povedať o reťazení, ktoré zahŕňa vyhľadávanie ukazovateľa v prepojenom zozname. Hašovacia tabuľka je jednoduché pole ukladajúce prvky KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Veľkosť tabuľky je mocnina dvoch, nie prvočíslo, pretože aplikácia masky pow2/AND vyžaduje jednu rýchlu inštrukciu, ale modulový operátor je oveľa pomalší. To je dôležité v prípade lineárneho snímania, pretože pri lineárnom vyhľadávaní tabuľky musí byť index slotu zabalený v každom slote. V dôsledku toho sa náklady na prevádzku pripočítajú modulo v každom slote.
Tabuľka ukladá iba kľúč a hodnotu pre každý prvok, nie hash kľúča. Keďže v tabuľke sú uložené iba 32-bitové kľúče, hash sa vypočíta veľmi rýchlo. Vyššie uvedený kód používa hash Murmur3, ktorý vykonáva iba niekoľko posunov, XOR a násobenia.
Hašovacia tabuľka používa techniky ochrany uzamknutia, ktoré sú nezávislé od poradia pamäte. Aj keď niektoré operácie zápisu narušia poradie iných takýchto operácií, tabuľka hash bude stále udržiavať správny stav. O tom si povieme nižšie. Táto technika funguje skvele s grafickými kartami, ktoré spúšťajú tisíce vlákien súčasne.
Kľúče a hodnoty v tabuľke hash sú inicializované tak, aby boli prázdne.
Kód je možné upraviť tak, aby spracoval aj 64-bitové kľúče a hodnoty. Kľúče vyžadujú atomické operácie čítania, zápisu a porovnávania a výmeny. A hodnoty vyžadujú atómové operácie čítania a zápisu. Našťastie v CUDA sú operácie čítania a zápisu pre 32- a 64-bitové hodnoty atómové, pokiaľ sú prirodzene zarovnané (pozri nižšie). ) a moderné grafické karty podporujú operácie 64-bitového atómového porovnávania a výmeny. Samozrejme, pri prechode na 64 bitov sa výkon mierne zníži.
Stav hash tabuľky
Každý pár kľúč – hodnota v hašovacej tabuľke môže mať jeden zo štyroch stavov:
- Kľúč a hodnota sú prázdne. V tomto stave je hašovacia tabuľka inicializovaná.
- Kľúč bol zapísaný, ale hodnota ešte nebola zapísaná. Ak iné vlákno práve číta údaje, vráti sa prázdne. To je normálne, to isté by sa stalo, keby iné vlákno vykonávania fungovalo o niečo skôr, a to hovoríme o súbežnej dátovej štruktúre.
- Zaznamenáva sa kľúč aj hodnota.
- Hodnota je dostupná pre iné vlákna vykonávania, ale kľúč ešte nie je. To sa môže stať, pretože programovací model CUDA má voľne usporiadaný model pamäte. To je normálne, v každom prípade je kľúč stále prázdny, aj keď hodnota už nie je.
Dôležitou nuansou je, že po zapísaní kľúča do slotu sa už nepohybuje - aj keď je kľúč vymazaný, o tom budeme hovoriť nižšie.
Kód hašovacej tabuľky dokonca funguje s voľne usporiadanými modelmi pamäte, v ktorých nie je známe poradie, v ktorom sa pamäť číta a zapisuje. Keď sa pozrieme na vkladanie, vyhľadávanie a odstraňovanie v hašovacej tabuľke, nezabudnite, že každý pár kľúč – hodnota je v jednom zo štyroch stavov opísaných vyššie.
Vkladanie do hašovacej tabuľky
Funkcia CUDA, ktorá vkladá páry kľúč – hodnota do hašovacej tabuľky, vyzerá takto:
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);
}
}
Ak chcete vložiť kľúč, kód iteruje pole hašovacej tabuľky počnúc hashom vloženého kľúča. Každý slot v poli vykoná operáciu atómového porovnania a výmeny, ktorá porovnáva kľúč v tomto slote s prázdnym. Ak sa zistí nesúlad, kľúč v slote sa aktualizuje vloženým kľúčom a potom sa vráti pôvodný kľúč slotu. Ak bol tento pôvodný kľúč prázdny alebo sa zhodoval s vloženým kľúčom, potom kód našiel vhodný slot na vloženie a vložil vloženú hodnotu do slotu.
Ak v jednom volaní jadra gpu_hashtable_insert() existuje viacero prvkov s rovnakým kľúčom, potom je možné do kľúča zapísať ktorúkoľvek z ich hodnôt. To sa považuje za normálne: jeden zo zápisov kľúč-hodnota počas volania bude úspešný, ale keďže sa to všetko deje paralelne v rámci niekoľkých vlákien vykonávania, nemôžeme predpovedať, ktorý zápis do pamäte bude posledný.
Vyhľadávanie v tabuľke hash
Kód pre vyhľadávanie kľúčov:
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);
}
}
Aby sme našli hodnotu kľúča uloženého v tabuľke, iterujeme cez pole počnúc hashom kľúča, ktorý hľadáme. V každom slote skontrolujeme, či je kľúč ten, ktorý hľadáme, a ak áno, vrátime jeho hodnotu. Tiež skontrolujeme, či je kľúč prázdny, a ak áno, vyhľadávanie prerušíme.
Ak kľúč nenájdeme, kód vráti prázdnu hodnotu.
Všetky tieto operácie vyhľadávania možno vykonávať súčasne prostredníctvom vkladania a vymazania. Každý pár v tabuľke bude mať jeden zo štyroch stavov opísaných vyššie pre tok.
Odstránenie v hašovacej tabuľke
Kód na vymazanie kľúčov:
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);
}
}
Vymazanie kľúča prebieha nezvyčajným spôsobom: kľúč necháme v tabuľke a jeho hodnotu (nie kľúč samotný) označíme ako prázdnu. Tento kód je veľmi podobný lookup(), okrem toho, že keď sa nájde zhoda na kľúči, jeho hodnota bude prázdna.
Ako je uvedené vyššie, akonáhle je kľúč zapísaný do slotu, už sa nepohybuje. Aj keď je prvok z tabuľky odstránený, kľúč zostáva na svojom mieste, jeho hodnota sa jednoducho vyprázdni. To znamená, že pre hodnotu slotu nemusíme použiť operáciu atomického zápisu, pretože nezáleží na tom, či je aktuálna hodnota prázdna alebo nie – aj tak sa vyprázdni.
Zmena veľkosti hašovacej tabuľky
Veľkosť hašovacej tabuľky môžete zmeniť tak, že vytvoríte väčšiu tabuľku a vložíte do nej neprázdne prvky zo starej tabuľky. Túto funkciu som neimplementoval, pretože som chcel zachovať jednoduchý vzorový kód. Navyše v programoch CUDA sa alokácia pamäte často vykonáva v hostiteľskom kóde, a nie v jadre CUDA.
Tento článok popisuje, ako upraviť takúto dátovú štruktúru chránenú zámkom.
konkurencieschopnosť
Vo vyššie uvedených útržkoch kódu funkcie gpu_hashtable_insert(), _lookup() и _delete() spracovať jeden pár kľúč – hodnota naraz. A nižšie gpu_hashtable_insert(), _lookup() и _delete() spracovať pole párov paralelne, každý pár v samostatnom vlákne vykonávania GPU:
// 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);
}
}
Hašovacia tabuľka odolná voči uzamknutiu podporuje súbežné vkladanie, vyhľadávanie a odstraňovanie. Pretože páry kľúč – hodnota sú vždy v jednom zo štyroch stavov a kľúče sa nepohybujú, tabuľka zaručuje správnosť aj pri súčasnom použití rôznych typov operácií.
Ak však spracujeme dávku vkladania a vymazania paralelne a ak vstupné pole párov obsahuje duplicitné kľúče, potom nebudeme schopní predpovedať, ktoré páry „vyhrajú“ – budú zapísané do hašovacej tabuľky ako posledné. Povedzme, že sme zavolali kód na vloženie so vstupným poľom párov A/0 B/1 A/2 C/3 A/4. Po dokončení kódu sa spáruje B/1 и C/3 sa v tabuľke zaručene vyskytnú, no zároveň sa v nej objaví ktorýkoľvek z párov A/0, A/2 alebo A/4. To môže a nemusí byť problém – všetko závisí od aplikácie. Možno vopred viete, že vo vstupnom poli nie sú žiadne duplicitné kľúče, alebo vám môže byť jedno, ktorá hodnota bola zapísaná ako posledná.
Ak je to pre vás problém, potom musíte rozdeliť duplicitné páry do rôznych systémových volaní CUDA. V CUDA sa každá operácia, ktorá volá jadro, vždy dokončí pred ďalším volaním jadra (aspoň v rámci jedného vlákna. V rôznych vláknach sa jadrá vykonávajú paralelne). Vo vyššie uvedenom príklade, ak zavoláte jedno jadro s A/0 B/1 A/2 C/3, a druhý s A/4, potom kľúč A dostane hodnotu 4.
Teraz si povedzme, či by funkcie mali lookup() и delete() použite obyčajný alebo nestály ukazovateľ na pole párov v hašovacej tabuľke. uvádza, že:
Kompilátor sa môže rozhodnúť optimalizovať čítanie a zápis do globálnej alebo zdieľanej pamäte... Tieto optimalizácie je možné vypnúť pomocou kľúčového slova
volatile: ... akýkoľvek odkaz na túto premennú je skompilovaný do skutočnej inštrukcie na čítanie alebo zápis z pamäte.
Úvahy o správnosti nevyžadujú aplikáciu volatile. Ak vykonávacie vlákno používa hodnotu uloženú vo vyrovnávacej pamäti z predchádzajúcej operácie čítania, bude používať mierne zastarané informácie. Ale stále je to informácia zo správneho stavu hašovacej tabuľky v určitom momente volania jadra. Ak potrebujete použiť najnovšie informácie, môžete použiť index volatile, ale potom sa výkon mierne zníži: podľa mojich testov sa pri vymazaní 32 miliónov prvkov rýchlosť znížila z 500 miliónov vymazaní/s na 450 miliónov vymazaní/s.
produktivita
V teste na vloženie 64 miliónov prvkov a vymazanie 32 miliónov z nich súťaž medzi std::unordered_map a pre GPU prakticky neexistuje hašovacia tabuľka:
std::unordered_map strávil 70 691 ms vkladaním a vyberaním prvkov a ich následným uvoľňovaním unordered_map (zbavenie sa miliónov prvkov si vyžaduje veľa času, pretože vo vnútri unordered_map vykoná sa viacnásobné pridelenie pamäte). úprimne povedané, std:unordered_map úplne iné obmedzenia. Ide o jedno vlákno CPU, ktoré podporuje kľúče a hodnoty akejkoľvek veľkosti, funguje dobre pri vysokej miere využitia a vykazuje stabilný výkon po viacnásobnom odstránení.
Trvanie hašovacej tabuľky pre GPU a medziprogramovú komunikáciu bolo 984 ms. To zahŕňa čas strávený umiestnením tabuľky do pamäte a jej vymazaním (jednorazové pridelenie 1 GB pamäte, čo v CUDA nejaký čas trvá), vkladaním a odstraňovaním prvkov a ich opakovaním. Do úvahy sa berú aj všetky kópie do az pamäte grafickej karty.
Dokončenie samotnej hašovacej tabuľky trvalo 271 ms. To zahŕňa čas, ktorý grafická karta strávi vkladaním a odstraňovaním prvkov, a neberie do úvahy čas strávený kopírovaním do pamäte a iterovaním výslednej tabuľky. Ak tabuľka GPU žije dlhú dobu alebo ak je hašovacia tabuľka obsiahnutá celá v pamäti grafickej karty (napríklad na vytvorenie hašovacej tabuľky, ktorú bude používať iný kód GPU a nie centrálny procesor), potom výsledok testu je relevantný.
Tabuľka hash pre grafickú kartu demonštruje vysoký výkon vďaka vysokej priepustnosti a aktívnej paralelizácii.
Obmedzenie
Architektúra hašovacej tabuľky má niekoľko problémov, o ktorých si treba uvedomiť:
- Lineárnemu snímaniu bráni zhlukovanie, čo spôsobuje, že kľúče v tabuľke nie sú umiestnené dokonale.
- Pomocou tejto funkcie sa kľúče neodstraňujú
deletea časom zaneriadia stôl.
Výsledkom je, že výkon hašovacej tabuľky sa môže postupne zhoršovať, najmä ak existuje dlhý čas a má veľa vložení a vymazaní. Jedným zo spôsobov, ako zmierniť tieto nevýhody, je prepracovať do novej tabuľky s pomerne nízkou mierou využitia a odfiltrovať odstránené kľúče počas prehadzovania.
Na ilustráciu popísaných problémov použijem vyššie uvedený kód na vytvorenie tabuľky so 128 miliónmi prvkov a prejdem cez 4 milióny prvkov, kým nezaplním 124 miliónov slotov (miera využitia približne 0,96). Tu je tabuľka výsledkov, každý riadok je volanie jadra CUDA na vloženie 4 miliónov nových prvkov do jednej hašovacej tabuľky:
Miera využitia
Trvanie vloženia 4 194 304 prvkov
0,00
11,608448 ms (361,314798 miliónov klávesov/s)
0,03
11,751424 ms (356,918799 miliónov klávesov/s)
0,06
11,942592 ms (351,205515 miliónov klávesov/s)
0,09
12,081120 ms (347,178429 miliónov klávesov/s)
0,12
12,242560 ms (342,600233 miliónov klávesov/s)
0,16
12,396448 ms (338,347235 miliónov klávesov/s)
0,19
12,533024 ms (334,660176 miliónov klávesov/s)
0,22
12,703328 ms (330,173626 miliónov klávesov/s)
0,25
12,884512 ms (325,530693 miliónov klávesov/s)
0,28
13,033472 ms (321,810182 miliónov klávesov/s)
0,31
13,239296 ms (316,807174 miliónov klávesov/s)
0,34
13,392448 ms (313,184256 miliónov klávesov/s)
0,37
13,624000 ms (307,861434 miliónov klávesov/s)
0,41
13,875520 ms (302,280855 miliónov klávesov/s)
0,44
14,126528 ms (296,909756 miliónov klávesov/s)
0,47
14,399328 ms (291,284699 miliónov klávesov/s)
0,50
14,690304 ms (285,515123 miliónov klávesov/s)
0,53
15,039136 ms (278,892623 miliónov klávesov/s)
0,56
15,478656 ms (270,973402 miliónov klávesov/s)
0,59
15,985664 ms (262,379092 miliónov klávesov/s)
0,62
16,668673 ms (251,627968 miliónov klávesov/s)
0,66
17,587200 ms (238,486174 miliónov klávesov/s)
0,69
18,690048 ms (224,413765 miliónov klávesov/s)
0,72
20,278816 ms (206,831789 miliónov klávesov/s)
0,75
22,545408 ms (186,038058 miliónov klávesov/s)
0,78
26,053312 ms (160,989275 miliónov klávesov/s)
0,81
31,895008 ms (131,503463 miliónov klávesov/s)
0,84
42,103294 ms (99,619378 miliónov klávesov/s)
0,87
61,849056 ms (67,815164 miliónov klávesov/s)
0,90
105,695999 ms (39,682713 miliónov klávesov/s)
0,94
240,204636 ms (17,461378 miliónov klávesov/s)
So zvyšujúcim sa využitím výkon klesá. Vo väčšine prípadov to nie je žiaduce. Ak aplikácia vloží prvky do tabuľky a následne ich zahodí (napríklad pri počítaní slov v knihe), tak to nie je problém. Ale ak aplikácia používa dlhotrvajúcu hašovaciu tabuľku (napríklad v grafickom editore na ukladanie neprázdnych častí obrázkov, kam používateľ často vkladá a odstraňuje informácie), potom môže byť toto správanie problematické.
A merali hĺbku sondovania hašovacej tabuľky po 64 miliónoch vložiek (faktor využitia 0,5). Priemerná hĺbka bola 0,4774, takže väčšina kláves bola buď v najlepšom možnom slote, alebo o jeden slot ďalej od najlepšej pozície. Maximálna hĺbka zvuku bola 60.
Potom som zmeral hĺbku sondovania na stole so 124 miliónmi vložiek (faktor využitia 0,97). Priemerná hĺbka už bola 10,1757 a maximálna - 6474 (!!). Výkon lineárneho snímania výrazne klesá pri vysokej miere využitia.
Najlepšie je udržiavať nízku mieru využitia tejto hašovacej tabuľky. Potom ale zvyšujeme výkon na úkor spotreby pamäte. Našťastie v prípade 32-bitových kľúčov a hodnôt sa to dá ospravedlniť. Ak vo vyššie uvedenom príklade v tabuľke so 128 miliónmi prvkov dodržíme koeficient využitia 0,25, potom do nej nemôžeme umiestniť viac ako 32 miliónov prvkov a zvyšných 96 miliónov slotov sa stratí – 8 bajtov na každý pár. , 768 MB stratenej pamäte.
Upozorňujeme, že hovoríme o strate pamäte grafickej karty, ktorá je cennejším zdrojom ako systémová pamäť. Hoci väčšina moderných desktopových grafických kariet, ktoré podporujú CUDA, má aspoň 4 GB pamäte (v čase písania článku má NVIDIA 2080 Ti 11 GB), stále by nebolo najmúdrejšie rozhodnutie o takéto sumy prísť.
Neskôr napíšem viac o vytváraní hašovacích tabuliek pre grafické karty, ktoré nemajú problémy s hĺbkou sondovania, ako aj o spôsoboch opätovného použitia vymazaných slotov.
Meranie hĺbky zvuku
Na určenie hĺbky skúmania kľúča môžeme extrahovať hash kľúča (jeho ideálny index do tabuľky) z jeho skutočného indexu tabuľky:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Kvôli kúzlu dvojkových dvojkových doplnkových binárnych čísel a skutočnosti, že kapacita hašovacej tabuľky je dva ku dvom, bude tento prístup fungovať aj vtedy, keď sa kľúčový index presunie na začiatok tabuľky. Zoberme si kľúč, ktorý zahašoval na 1, ale je vložený do slotu 3. Potom pre stôl s kapacitou 4 dostaneme (3 — 1) & 3, čo je ekvivalentné 2.
Záver
Ak máte otázky alebo pripomienky, pošlite mi e-mail na adresu alebo otvorte novú tému v.
Tento kód bol napísaný na základe inšpirácie z vynikajúcich článkov:
V budúcnosti budem naďalej písať o implementáciách hashovacích tabuliek pre grafické karty a analyzovať ich výkon. Moje plány zahŕňajú reťazenie, Robin Hood hašovanie a kukučkové hašovanie pomocou atómových operácií v dátových štruktúrach, ktoré sú priateľské k GPU.
Zdroj: hab.com