Zveřejnil jsem to na Github .
Je to jednoduchá hashovací tabulka GPU schopná zpracovat stovky milionů vložení za sekundu. Na mém notebooku NVIDIA GTX 1060 kód vloží 64 milionů náhodně vygenerovaných párů klíč-hodnota za přibližně 210 ms a odstraní 32 milionů párů za přibližně 64 ms.
To znamená, že rychlost na notebooku je přibližně 300 milionů vložení/s a 500 milionů smazání/s.
Tabulka je napsána v CUDA, i když stejnou techniku lze aplikovat na HLSL nebo GLSL. Implementace má několik omezení pro zajištění vysokého výkonu na grafické kartě:
- Zpracovávají se pouze 32bitové klíče a stejné hodnoty.
- Hašovací tabulka má pevnou velikost.
- A tato velikost se musí rovnat mocnině dvěma.
Pro klíče a hodnoty si musíte rezervovat jednoduchý oddělovací znak (ve výše uvedeném kódu je to 0xffffffff).
Hash tabulka bez zámků
Hashovací tabulka používá otevřené adresování s , to znamená, že je to jednoduše pole párů klíč-hodnota, které je uloženo v paměti a má vynikající výkon mezipaměti. Totéž nelze říci o řetězení, které zahrnuje hledání ukazatele v propojeném seznamu. Hašovací tabulka je jednoduché pole ukládající prvky KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Velikost tabulky je mocnina dvou, ne prvočíslo, protože k aplikaci masky pow2/AND stačí jedna rychlá instrukce, ale modulový operátor je mnohem pomalejší. To je důležité v případě lineárního snímání, protože při lineárním vyhledávání v tabulce musí být index slotu zabalen do každého slotu. A v důsledku toho se náklady na provoz přidávají modulo v každém slotu.
Tabulka ukládá pouze klíč a hodnotu pro každý prvek, nikoli hash klíče. Protože tabulka ukládá pouze 32bitové klíče, hash se vypočítá velmi rychle. Výše uvedený kód používá hash Murmur3, který provádí pouze několik posunů, XOR a násobení.
Tabulka hash používá techniky ochrany zamykání, které jsou nezávislé na pořadí paměti. I když některé operace zápisu naruší pořadí jiných takových operací, hashovací tabulka bude stále udržovat správný stav. O tom si povíme níže. Tato technika funguje skvěle s grafickými kartami, které běží na tisících vláken současně.
Klíče a hodnoty v hash tabulce jsou inicializovány tak, aby byly prázdné.
Kód lze upravit tak, aby zpracovával i 64bitové klíče a hodnoty. Klíče vyžadují atomické operace čtení, zápisu a porovnání a výměny. A hodnoty vyžadují atomické operace čtení a zápisu. Naštěstí v CUDA jsou operace čtení a zápisu pro 32bitové a 64bitové hodnoty atomické, pokud jsou přirozeně zarovnány (viz níže). ) a moderní grafické karty podporují 64bitové atomové operace porovnávání a výměny. Při přechodu na 64 bitů se výkon samozřejmě mírně sníží.
Stav hashovací tabulky
Každý pár klíč–hodnota v hašovací tabulce může mít jeden ze čtyř stavů:
- Klíč a hodnota jsou prázdné. V tomto stavu je hashovací tabulka inicializována.
- Klíč byl zapsán, ale hodnota ještě nebyla zapsána. Pokud jiné vlákno aktuálně čte data, vrátí se prázdné. To je normální, totéž by se stalo, kdyby jiné vlákno spouštění fungovalo o něco dříve, a to mluvíme o souběžné datové struktuře.
- Zaznamenává se klíč i hodnota.
- Hodnota je dostupná pro ostatní vlákna provádění, ale klíč zatím není. To se může stát, protože programovací model CUDA má volně uspořádaný model paměti. To je normální, v každém případě je klíč stále prázdný, i když hodnota již není.
Důležitou nuancí je, že jakmile je klíč zapsán do slotu, již se nepohybuje - i když je klíč odstraněn, o tom budeme hovořit níže.
Kód hashovací tabulky funguje dokonce i s volně uspořádanými modely paměti, ve kterých je neznámé pořadí, ve kterém je paměť čtena a zapisována. Když se podíváme na vkládání, vyhledávání a mazání v hašovací tabulce, nezapomeňte, že každý pár klíč-hodnota je v jednom ze čtyř výše popsaných stavů.
Vkládání do hash tabulky
Funkce CUDA, která vkládá páry klíč–hodnota do hašovací tabulky, vypadá 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);
}
}
Chcete-li vložit klíč, kód iteruje polem hash tabulky počínaje hashem vloženého klíče. Každý slot v poli provede operaci atomického porovnání a výměny, která porovná klíč v tomto slotu s prázdným. Pokud je zjištěna neshoda, klíč ve slotu je aktualizován vloženým klíčem a poté je vrácen původní klíč slotu. Pokud byl tento původní klíč prázdný nebo se shodoval s vloženým klíčem, pak kód našel vhodný slot pro vložení a vložil vloženou hodnotu do slotu.
Pokud v jednom volání jádra gpu_hashtable_insert() existuje více prvků se stejným klíčem, pak lze do slotu klíče zapsat kteroukoli z jejich hodnot. To je považováno za normální: jeden ze zápisů klíč–hodnota během volání bude úspěšný, ale protože se to vše děje paralelně v několika vláknech provádění, nemůžeme předvídat, který zápis do paměti bude poslední.
Vyhledávání v hash tabulce
Kód pro vyhledávání klíčů:
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);
}
}
Abychom našli hodnotu klíče uloženého v tabulce, iterujeme polem počínaje hash klíče, který hledáme. V každém slotu zkontrolujeme, zda je klíč ten, který hledáme, a pokud ano, vrátíme jeho hodnotu. Zkontrolujeme také, zda je klíč prázdný, a pokud ano, vyhledávání přerušíme.
Pokud klíč nenajdeme, vrátí kód prázdnou hodnotu.
Všechny tyto vyhledávací operace lze provádět současně prostřednictvím vkládání a mazání. Každý pár v tabulce bude mít jeden ze čtyř stavů popsaných výše pro tok.
Mazání v hašovací tabulce
Kód pro smazání klíčů:
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);
}
}
Smazání klíče se provádí neobvyklým způsobem: klíč necháme v tabulce a jeho hodnotu (nikoli klíč samotný) označíme jako prázdnou. Tento kód je velmi podobný lookup(), kromě toho, že když je nalezena shoda na klíči, jeho hodnota bude prázdná.
Jak bylo uvedeno výše, jakmile je klíč zapsán do slotu, již se nepohybuje. I když je prvek z tabulky odstraněn, klíč zůstává na svém místě, jeho hodnota se jednoduše vyprázdní. To znamená, že pro hodnotu slotu nemusíme používat operaci atomického zápisu, protože nezáleží na tom, zda je aktuální hodnota prázdná nebo ne – i tak bude prázdná.
Změna velikosti hash tabulky
Velikost hash tabulky můžete změnit vytvořením větší tabulky a vložením neprázdných prvků ze staré tabulky do ní. Tuto funkci jsem neimplementoval, protože jsem chtěl zachovat ukázkový kód jednoduchý. Navíc v programech CUDA se alokace paměti často provádí v hostitelském kódu spíše než v jádře CUDA.
Tento článek popisuje, jak upravit takovou datovou strukturu chráněnou zámkem.
Konkurenceschopnost
Ve výše uvedených fragmentech kódu funkce gpu_hashtable_insert(), _lookup() и _delete() zpracovávat vždy jeden pár klíč–hodnota. A nižší gpu_hashtable_insert(), _lookup() и _delete() zpracujte paralelně pole párů, každý pár v samostatném vláknu provádění 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šovací tabulka odolná proti uzamčení podporuje souběžné vkládání, vyhledávání a mazání. Protože páry klíč-hodnota jsou vždy v jednom ze čtyř stavů a klíče se nepohybují, tabulka zaručuje správnost i při současném použití různých typů operací.
Pokud však zpracováváme dávku vkládání a mazání paralelně a pokud vstupní pole párů obsahuje duplicitní klíče, pak nebudeme schopni předpovědět, které páry „vyhrají“ – budou zapsány do hašovací tabulky jako poslední. Řekněme, že jsme zavolali vkládací kód se vstupním polem párů A/0 B/1 A/2 C/3 A/4. Po dokončení kódu dojde k párování B/1 и C/3 jsou zaručeně přítomni v tabulce, ale zároveň se v ní objeví kterýkoli z párů A/0, A/2 nebo A/4. To může, ale nemusí být problém – vše závisí na aplikaci. Možná předem víte, že ve vstupním poli nejsou žádné duplicitní klíče, nebo vám může být jedno, která hodnota byla zapsána jako poslední.
Pokud je to pro vás problém, musíte duplicitní páry rozdělit do různých systémových volání CUDA. V CUDA se každá operace, která volá jádro, vždy dokončí před dalším voláním jádra (alespoň v rámci jednoho vlákna. V různých vláknech se jádra spouštějí paralelně). Ve výše uvedeném příkladu, pokud zavoláte jedno jádro s A/0 B/1 A/2 C/3, a druhý s A/4, pak klíč A získá hodnotu 4.
Nyní pojďme mluvit o tom, zda by funkce měly lookup() и delete() použijte prostý nebo nestálý ukazatel na pole párů v tabulce hash. Tvrdí, že:
Kompilátor se může rozhodnout optimalizovat čtení a zápis do globální nebo sdílené paměti... Tyto optimalizace lze zakázat pomocí klíčového slova
volatile: ... jakýkoli odkaz na tuto proměnnou je zkompilován do skutečné instrukce pro čtení nebo zápis do paměti.
Úvahy o správnosti nevyžadují aplikaci volatile. Pokud spouštěcí vlákno používá hodnotu uloženou v mezipaměti z dřívější operace čtení, bude používat mírně zastaralé informace. Ale přesto je to informace ze správného stavu hash tabulky v určitém okamžiku volání jádra. Pokud potřebujete použít nejnovější informace, můžete použít rejstřík volatile, ale pak se výkon mírně sníží: podle mých testů se při smazání 32 milionů prvků snížila rychlost z 500 milionů smazání/s na 450 milionů smazání/s.
Производительность
V testu na vložení 64 milionů prvků a vymazání 32 milionů z nich soutěž mezi std::unordered_map a pro GPU prakticky neexistuje žádná hashovací tabulka:
std::unordered_map strávil 70 691 ms vkládáním a vyjímáním prvků a jejich následným uvolňováním unordered_map (zbavení se milionů prvků zabere spoustu času, protože uvnitř unordered_map je provedeno více alokací paměti). upřímně řečeno, std:unordered_map úplně jiná omezení. Jedná se o jediné vlákno CPU, které se spouští, podporuje páry klíč–hodnota jakékoli velikosti, funguje dobře při vysokých rychlostech využití a vykazuje stabilní výkon po vícenásobných smazáních.
Doba trvání hashovací tabulky pro GPU a meziprogramovou komunikaci byla 984 ms. To zahrnuje čas strávený umístěním tabulky do paměti a jejím mazáním (jednorázové přidělení 1 GB paměti, což v CUDA nějakou dobu zabere), vkládáním a mazáním prvků a jejich opakováním. Berou se v úvahu také všechny kopie do az paměti grafické karty.
Dokončení samotné hashovací tabulky trvalo 271 ms. To zahrnuje čas strávený grafickou kartou vkládáním a mazáním prvků a nebere v úvahu čas strávený kopírováním do paměti a iterací výsledné tabulky. Pokud tabulka GPU žije dlouhou dobu nebo pokud je hashovací tabulka obsažena celá v paměti grafické karty (například pro vytvoření hashovací tabulky, kterou bude používat jiný kód GPU a ne centrální procesor), pak výsledek testu je relevantní.
Hashovací tabulka pro grafickou kartu ukazuje vysoký výkon díky vysoké propustnosti a aktivní paralelizaci.
Omezení
Architektura hashovací tabulky má několik problémů, o kterých je třeba si uvědomit:
- Lineární sondování je ztíženo shlukováním, což způsobuje, že klíče v tabulce nejsou umístěny dokonale.
- Pomocí této funkce se klíče neodstraňují
deletea časem zaneřádí stůl.
Výsledkem je, že výkon hashovací tabulky se může postupně zhoršovat, zejména pokud existuje po dlouhou dobu a má mnoho vložení a odstranění. Jedním ze způsobů, jak tyto nevýhody zmírnit, je předělat do nové tabulky s poměrně nízkou mírou využití a odfiltrovat odstraněné klíče během přehánění.
Pro ilustraci popsaných problémů použiji výše uvedený kód k vytvoření tabulky se 128 miliony prvků a budu procházet 4 miliony prvků, dokud nezaplním 124 milionů slotů (míra využití asi 0,96). Zde je výsledková tabulka, každý řádek je voláním jádra CUDA pro vložení 4 milionů nových prvků do jedné hashovací tabulky:
Míra využití
Délka vložení 4 194 304 prvků
0,00
11,608448 ms (361,314798 milionů klíčů/s)
0,03
11,751424 ms (356,918799 milionů klíčů/s)
0,06
11,942592 ms (351,205515 milionů klíčů/s)
0,09
12,081120 ms (347,178429 milionů klíčů/s)
0,12
12,242560 ms (342,600233 milionů klíčů/s)
0,16
12,396448 ms (338,347235 milionů klíčů/s)
0,19
12,533024 ms (334,660176 milionů klíčů/s)
0,22
12,703328 ms (330,173626 milionů klíčů/s)
0,25
12,884512 ms (325,530693 milionů klíčů/s)
0,28
13,033472 ms (321,810182 milionů klíčů/s)
0,31
13,239296 ms (316,807174 milionů klíčů/s)
0,34
13,392448 ms (313,184256 milionů klíčů/s)
0,37
13,624000 ms (307,861434 milionů klíčů/s)
0,41
13,875520 ms (302,280855 milionů klíčů/s)
0,44
14,126528 ms (296,909756 milionů klíčů/s)
0,47
14,399328 ms (291,284699 milionů klíčů/s)
0,50
14,690304 ms (285,515123 milionů klíčů/s)
0,53
15,039136 ms (278,892623 milionů klíčů/s)
0,56
15,478656 ms (270,973402 milionů klíčů/s)
0,59
15,985664 ms (262,379092 milionů klíčů/s)
0,62
16,668673 ms (251,627968 milionů klíčů/s)
0,66
17,587200 ms (238,486174 milionů klíčů/s)
0,69
18,690048 ms (224,413765 milionů klíčů/s)
0,72
20,278816 ms (206,831789 milionů klíčů/s)
0,75
22,545408 ms (186,038058 milionů klíčů/s)
0,78
26,053312 ms (160,989275 milionů klíčů/s)
0,81
31,895008 ms (131,503463 milionů klíčů/s)
0,84
42,103294 ms (99,619378 milionů klíčů/s)
0,87
61,849056 ms (67,815164 milionů klíčů/s)
0,90
105,695999 ms (39,682713 milionů klíčů/s)
0,94
240,204636 ms (17,461378 milionů klíčů/s)
S rostoucím využitím se výkon snižuje. To ve většině případů není žádoucí. Pokud aplikace vloží prvky do tabulky a následně je zahodí (například při počítání slov v knize), pak to není problém. Pokud ale aplikace používá dlouhodobou hashovací tabulku (například v grafickém editoru k ukládání neprázdných částí obrázků, kam uživatel často vkládá a odstraňuje informace), pak může být toto chování problematické.
A změřil hloubku sondování hash tabulky po 64 milionech vložek (faktor využití 0,5). Průměrná hloubka byla 0,4774, takže většina kláves byla buď v nejlepším možném slotu, nebo jeden slot od nejlepší pozice. Maximální hloubka zvuku byla 60.
Poté jsem změřil hloubku sondování na stole se 124 miliony břitových destiček (faktor využití 0,97). Průměrná hloubka byla již 10,1757 a maximální - 6474 (!!). Výkon lineárního snímání výrazně klesá při vysoké míře využití.
Nejlepší je udržovat nízkou míru využití této hashovací tabulky. Pak ale zvyšujeme výkon na úkor spotřeby paměti. Naštěstí v případě 32bitových klíčů a hodnot to lze ospravedlnit. Pokud ve výše uvedeném příkladu v tabulce se 128 miliony prvků zachováme faktor využití 0,25, pak do ní nemůžeme umístit více než 32 milionů prvků a zbývajících 96 milionů slotů bude ztraceno - 8 bajtů pro každý pár , 768 MB ztracené paměti.
Vezměte prosím na vědomí, že mluvíme o ztrátě paměti grafické karty, což je cennější zdroj než systémová paměť. Většina moderních grafických karet pro stolní počítače, které podporují CUDA, má sice alespoň 4 GB paměti (v době psaní tohoto článku má NVIDIA 2080 Ti 11 GB), přesto by nebylo nejmoudřejší rozhodnutí o takové částky přijít.
Později napíšu více o vytváření hashovacích tabulek pro grafické karty, které nemají problémy s hloubkou sondování, a také o způsobech opětovného použití smazaných slotů.
Měření hloubky ozvučení
Abychom určili hloubku zkoumání klíče, můžeme extrahovat hash klíče (jeho ideální index tabulky) z jeho skutečného indexu tabulky:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Kvůli kouzlu dvojkových dvojkových doplňkových binárních čísel a skutečnosti, že kapacita hašovací tabulky je dvě ku dvojce, bude tento přístup fungovat, i když se index klíče přesune na začátek tabulky. Vezměme klíč, který zahašoval na 1, ale je vložen do slotu 3. Pak pro stůl s kapacitou 4 dostaneme (3 — 1) & 3, což se rovná 2.
Závěr
Máte-li dotazy nebo připomínky, napište mi na e-mail nebo otevřít nové téma v.
Tento kód byl napsán na základě inspirace z vynikajících článků:
V budoucnu budu nadále psát o implementacích hashovacích tabulek pro grafické karty a analyzovat jejich výkon. Moje plány zahrnují řetězení, Robin Hood hašování a kukaččí hašování pomocí atomických operací v datových strukturách, které jsou přátelské k GPU.
Zdroj: www.habr.com