Objavio sam ga na Github-u .
To je jednostavna GPU hash tablica koja može obraditi stotine miliona umetanja u sekundi. Na mom NVIDIA GTX 1060 laptopu, kod ubacuje 64 miliona nasumično generiranih parova ključ/vrijednost za oko 210 ms i uklanja 32 miliona parova za oko 64 ms.
Odnosno, brzina na laptopu je približno 300 miliona umetanja u sekundi i 500 miliona brisanja u sekundi.
Tabela je napisana u CUDA-i, iako se ista tehnika može primijeniti na HLSL ili GLSL. Implementacija ima nekoliko ograničenja kako bi se osigurale visoke performanse na video kartici:
- Obrađuju se samo 32-bitni ključevi i iste vrijednosti.
- Haš tabela ima fiksnu veličinu.
- I ova veličina mora biti jednaka dva na potenciju.
Za ključeve i vrijednosti, morate rezervirati jednostavan graničnik (u gornjem kodu to je 0xffffffff).
Hash tablica bez brava
Haš tabela koristi otvoreno adresiranje sa , to jest, to je jednostavno niz parova ključ/vrijednost koji je pohranjen u memoriji i ima superiorne performanse keš memorije. Isto se ne može reći za ulančavanje, koje uključuje traženje pokazivača na povezanoj listi. Haš tabela je jednostavan niz za pohranjivanje elemenata KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Veličina tabele je stepen dva, a ne prost broj, jer je jedna brza instrukcija dovoljna za primenu pow2/AND maske, ali modul operator je mnogo sporiji. Ovo je važno u slučaju linearnog sondiranja, jer u linearnom traženju u tabeli indeks slota mora biti umotan u svaki slot. I kao rezultat toga, cijena operacije se dodaje modulom u svaki slot.
Tabela pohranjuje samo ključ i vrijednost za svaki element, a ne heš ključa. Pošto tabela pohranjuje samo 32-bitne ključeve, heš se izračunava vrlo brzo. Kod iznad koristi hash Murmur3, koji izvodi samo nekoliko pomaka, XOR-ova i množenja.
Haš tablica koristi tehnike zaštite zaključavanja koje su neovisne o redoslijedu memorije. Čak i ako neke operacije pisanja poremete redoslijed drugih takvih operacija, hash tablica će i dalje održavati ispravno stanje. O tome ćemo razgovarati u nastavku. Tehnika odlično funkcionira s video karticama koje istovremeno pokreću hiljade niti.
Ključevi i vrijednosti u hash tablici su inicijalizirani na prazne.
Kod se može modificirati i za rukovanje 64-bitnim ključevima i vrijednostima. Ključevi zahtijevaju atomske operacije čitanja, pisanja i poređenja i zamjene. A vrijednosti zahtijevaju atomske operacije čitanja i pisanja. Na sreću, u CUDA-i, operacije čitanja i pisanja za 32- i 64-bitne vrijednosti su atomske sve dok su prirodno usklađene (pogledajte dolje). ), a moderne video kartice podržavaju 64-bitne atomske operacije usporedbe i razmjene. Naravno, kada se pređe na 64 bita, performanse će se malo smanjiti.
Hash stanje tabele
Svaki par ključ/vrijednost u hash tabeli može imati jedno od četiri stanja:
- Ključ i vrijednost su prazni. U ovom stanju, heš tablica je inicijalizirana.
- Ključ je zapisan, ali vrijednost još nije zapisana. Ako druga nit trenutno čita podatke, onda se vraća prazno. Ovo je normalno, ista stvar bi se desila da je druga nit izvršenja proradila malo ranije, a govorimo o konkurentnoj strukturi podataka.
- Snimaju se i ključ i vrijednost.
- Vrijednost je dostupna drugim nitima izvršenja, ali ključ još nije. Ovo se može dogoditi zato što model programiranja CUDA ima labavo uređen memorijski model. Ovo je normalno; u svakom slučaju, ključ je i dalje prazan, čak i ako vrijednost više nije takva.
Važna nijansa je da nakon što je ključ upisan u slot, više se ne pomiče - čak i ako je ključ izbrisan, o tome ćemo govoriti u nastavku.
Kod hash table radi čak i sa labavo uređenim memorijskim modelima u kojima je redosled čitanja i upisivanja memorije nepoznat. Dok gledamo umetanje, traženje i brisanje u hash tabeli, zapamtite da je svaki par ključ/vrijednost u jednom od četiri gore opisana stanja.
Umetanje u hash tabelu
Funkcija CUDA koja ubacuje parove ključ/vrijednost u hash tablicu izgleda ovako:
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);
}
}
Da bi umetnuo ključ, kod se ponavlja kroz niz tablice heš počevši od hash umetnutog ključa. Svaki slot u nizu izvodi atomsku operaciju usporedbe i zamjene koja uspoređuje ključ u tom slotu sa praznim. Ako se otkrije nepodudaranje, ključ u utoru se ažurira umetnutim ključem, a zatim se vraća originalni ključ utora. Ako je ovaj originalni ključ bio prazan ili je odgovarao umetnutom ključu, tada je kod pronašao odgovarajući utor za umetanje i umetnuo umetnutu vrijednost u utor.
Ako u jednom pozivu kernela gpu_hashtable_insert() postoji više elemenata s istim ključem, tada se bilo koja od njihovih vrijednosti može upisati u utor za ključ. Ovo se smatra normalnim: jedno od upisa ključ/vrijednost tokom poziva će uspjeti, ali pošto se sve ovo dešava paralelno unutar nekoliko niti izvršavanja, ne možemo predvidjeti koje će upisivanje u memoriju biti posljednje.
Traženje hash tablice
Kod za traženje ključeva:
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);
}
}
Da bismo pronašli vrijednost ključa pohranjenog u tabeli, ponavljamo niz niz počevši od hash ključa koji tražimo. U svakom slotu provjeravamo da li je ključ onaj koji tražimo i ako jeste, vraćamo njegovu vrijednost. Također provjeravamo da li je ključ prazan i ako jeste, prekidamo pretragu.
Ako ne možemo pronaći ključ, kod vraća praznu vrijednost.
Sve ove operacije pretraživanja mogu se izvoditi istovremeno umetanjem i brisanjem. Svaki par u tabeli će imati jedno od četiri gore opisana stanja za tok.
Brisanje u hash tabeli
Kod za brisanje ključeva:
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);
}
}
Brisanje ključa se vrši na neobičan način: ključ ostavljamo u tabeli i njegovu vrijednost (ne sam ključ) označavamo kao praznu. Ovaj kod je vrlo sličan lookup(), osim što kada se pronađe podudaranje na ključu, čini njegovu vrijednost praznom.
Kao što je gore pomenuto, kada se ključ upiše u slot, on se više ne pomera. Čak i kada se element izbriše iz tabele, ključ ostaje na mestu, njegova vrednost jednostavno postaje prazna. To znači da ne trebamo koristiti atomsku operaciju pisanja za vrijednost slota, jer nije bitno da li je trenutna vrijednost prazna ili ne - ona će i dalje postati prazna.
Promjena veličine hash tablice
Možete promijeniti veličinu hash tablice kreiranjem veće tablice i umetanjem nepraznih elemenata iz stare tablice u nju. Nisam implementirao ovu funkcionalnost jer sam želio da uzorak koda bude jednostavan. Štaviše, u CUDA programima, alokacija memorije se često vrši u host kodu, a ne u CUDA kernelu.
Clanak opisuje kako modificirati takvu strukturu podataka zaštićenu zaključavanjem.
Konkurentnost
U gornjim isječcima koda funkcije gpu_hashtable_insert(), _lookup() и _delete() obraditi jedan po jedan par ključ/vrijednost. I niže gpu_hashtable_insert(), _lookup() и _delete() obraditi niz parova paralelno, svaki par u zasebnoj GPU izvršnoj niti:
// 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š tabela otporna na zaključavanje podržava istovremene umetanja, pretraživanja i brisanja. Budući da su parovi ključ/vrijednost uvijek u jednom od četiri stanja i ključevi se ne pomiču, tabela garantuje ispravnost čak i kada se različite vrste operacija koriste istovremeno.
Međutim, ako paralelno obrađujemo seriju umetanja i brisanja i ako ulazni niz parova sadrži duple ključeve, tada nećemo moći predvidjeti koji će parovi „pobijediti“ – bit će zadnji upisani u hash tablicu. Recimo da smo pozvali kod za umetanje sa ulaznim nizom parova A/0 B/1 A/2 C/3 A/4. Kada se kod završi, uparite B/1 и C/3 garantovano će biti prisutni u tabeli, ali će se istovremeno u njoj pojaviti bilo koji od parova A/0, A/2 ili A/4. Ovo može biti problem, ali i ne mora - sve ovisi o aplikaciji. Možda znate unaprijed da u nizu unosa nema duplih ključeva, ili vam možda nije važno koja je vrijednost zadnja upisana.
Ako vam ovo predstavlja problem, onda morate odvojiti duple parove u različite CUDA sistemske pozive. U CUDA-i, svaka operacija koja poziva kernel uvijek se završava prije sljedećeg poziva kernela (barem unutar jedne niti. U različitim nitima, kerneli se izvršavaju paralelno). U gornjem primjeru, ako pozovete jedno jezgro sa A/0 B/1 A/2 C/3, a drugi sa A/4, zatim ključ A dobiće vrednost 4.
Hajde sada da razgovaramo o tome treba li funkcije lookup() и delete() koristite običan ili promjenjiv pokazivač na niz parova u hash tablici. navodi da:
Kompajler može izabrati da optimizuje čitanje i upisivanje u globalnu ili zajedničku memoriju... Ove optimizacije se mogu onemogućiti pomoću ključne reči
volatile: ... svaka referenca na ovu varijablu kompajlira se u instrukciju za čitanje ili pisanje stvarne memorije.
Razmatranja ispravnosti ne zahtijevaju primjenu volatile. Ako izvršna nit koristi keširanu vrijednost iz prethodne operacije čitanja, tada će koristiti malo zastarjele informacije. Ali ipak, ovo je informacija iz ispravnog stanja hash tablice u određenom trenutku poziva kernela. Ako trebate koristiti najnovije informacije, možete koristiti indeks volatile, ali tada će se performanse neznatno smanjiti: prema mojim testovima, prilikom brisanja 32 miliona elemenata, brzina se smanjila sa 500 miliona brisanja u sekundi na 450 miliona brisanja u sekundi.
Produktivnost
U testu za umetanje 64 miliona elemenata i brisanje njih 32 miliona, konkurencija između std::unordered_map i praktično ne postoji hash tablica za GPU:
std::unordered_map proveo 70 ms umetajući i uklanjajući elemente i potom ih oslobađajući unordered_map (oslobađanje od miliona elemenata oduzima dosta vremena, jer unutra unordered_map vrši se višestruka dodjela memorije). Iskreno govoreći, std:unordered_map potpuno drugačija ograničenja. To je jedna CPU nit izvršenja, podržava ključ/vrijednost bilo koje veličine, dobro radi pri visokim stopama iskorištenja i pokazuje stabilne performanse nakon višestrukih brisanja.
Trajanje hash tablice za GPU i međuprogramsku komunikaciju je bilo 984 ms. Ovo uključuje vrijeme utrošeno na postavljanje tablice u memoriju i njeno brisanje (jednokratno dodjeljivanje 1 GB memorije, što traje neko vrijeme u CUDA-i), umetanje i brisanje elemenata i ponavljanje preko njih. Sve kopije u i iz memorije video kartice se također uzimaju u obzir.
Sama heš tabela se završila za 271 ms. Ovo uključuje vrijeme koje je video kartica potrošila na umetanje i brisanje elemenata, i ne uzima u obzir vrijeme utrošeno na kopiranje u memoriju i iteraciju preko rezultirajuće tablice. Ako GPU tabela živi dugo vremena, ili ako je hash tabela u potpunosti sadržana u memoriji video kartice (na primjer, za kreiranje hash tablice koju će koristiti drugi GPU kod, a ne centralni procesor), tada rezultat testa je relevantan.
Heš tablica za video karticu pokazuje visoke performanse zbog velike propusnosti i aktivne paralelizacije.
mane
Arhitektura hash tablice ima nekoliko problema kojih treba biti svjestan:
- Linearno ispitivanje je otežano grupiranjem, što uzrokuje da ključevi u tabeli budu postavljeni manje nego savršeno.
- Tipke se ne uklanjaju korištenjem funkcije
deletei vremenom zatrpaju sto.
Kao rezultat toga, performanse hash tablice mogu postepeno degradirati, posebno ako postoji dugo vremena i ima brojna umetanja i brisanja. Jedan od načina da se ublaže ovi nedostaci je rehasiranje u novu tabelu sa prilično niskom stopom iskorištenja i filtriranje uklonjenih ključeva tokom ponovnog haširanja.
Da bih ilustrovao opisane probleme, koristiću gornji kod za kreiranje tabele sa 128 miliona elemenata i petlju kroz 4 miliona elemenata dok ne popunim 124 miliona mesta (stopa iskorišćenja od oko 0,96). Evo tabele rezultata, svaki red je poziv CUDA kernela za umetanje 4 miliona novih elemenata u jednu hash tabelu:
Stopa iskorištenja
Trajanje umetanja 4 elementa
0,00
11,608448 ms (361,314798 miliona ključeva/sek.)
0,03
11,751424 ms (356,918799 miliona ključeva/sek.)
0,06
11,942592 ms (351,205515 miliona ključeva/sek.)
0,09
12,081120 ms (347,178429 miliona ključeva/sek.)
0,12
12,242560 ms (342,600233 miliona ključeva/sek.)
0,16
12,396448 ms (338,347235 miliona ključeva/sek.)
0,19
12,533024 ms (334,660176 miliona ključeva/sek.)
0,22
12,703328 ms (330,173626 miliona ključeva/sek.)
0,25
12,884512 ms (325,530693 miliona ključeva/sek.)
0,28
13,033472 ms (321,810182 miliona ključeva/sek.)
0,31
13,239296 ms (316,807174 miliona ključeva/sek.)
0,34
13,392448 ms (313,184256 miliona ključeva/sek.)
0,37
13,624000 ms (307,861434 miliona ključeva/sek.)
0,41
13,875520 ms (302,280855 miliona ključeva/sek.)
0,44
14,126528 ms (296,909756 miliona ključeva/sek.)
0,47
14,399328 ms (291,284699 miliona ključeva/sek.)
0,50
14,690304 ms (285,515123 miliona ključeva/sek.)
0,53
15,039136 ms (278,892623 miliona ključeva/sek.)
0,56
15,478656 ms (270,973402 miliona ključeva/sek.)
0,59
15,985664 ms (262,379092 miliona ključeva/sek.)
0,62
16,668673 ms (251,627968 miliona ključeva/sek.)
0,66
17,587200 ms (238,486174 miliona ključeva/sek.)
0,69
18,690048 ms (224,413765 miliona ključeva/sek.)
0,72
20,278816 ms (206,831789 miliona ključeva/sek.)
0,75
22,545408 ms (186,038058 miliona ključeva/sek.)
0,78
26,053312 ms (160,989275 miliona ključeva/sek.)
0,81
31,895008 ms (131,503463 miliona ključeva/sek.)
0,84
42,103294 ms (99,619378 miliona ključeva/sek.)
0,87
61,849056 ms (67,815164 miliona ključeva/sek.)
0,90
105,695999 ms (39,682713 miliona ključeva/sek.)
0,94
240,204636 ms (17,461378 miliona ključeva/sek.)
Kako se korištenje povećava, performanse se smanjuju. Ovo u većini slučajeva nije poželjno. Ako aplikacija ubacuje elemente u tabelu, a zatim ih odbacuje (na primjer, kada broji riječi u knjizi), onda to nije problem. Ali ako aplikacija koristi dugovječnu hash tablicu (na primjer, u grafičkom uređivaču za pohranjivanje nepraznih dijelova slika gdje korisnik često ubacuje i briše informacije), onda ovo ponašanje može biti problematično.
I izmerio dubinu ispitivanja heš tabele nakon 64 miliona umetanja (faktor iskorišćenja 0,5). Prosječna dubina je bila 0,4774, tako da je većina ključeva bila ili u najboljem mogućem slotu ili jedan slot udaljena od najbolje pozicije. Maksimalna dubina sondiranja bila je 60.
Zatim sam izmjerio dubinu sondiranja na stolu sa 124 miliona umetaka (faktor iskorištenja 0,97). Prosječna dubina je već bila 10,1757, a maksimalna - 6474 (!!). Performanse linearnog sensinga značajno opadaju pri visokim stopama iskorištenja.
Najbolje je držati nisku stopu iskorištenja ove hash tablice. Ali onda povećavamo performanse na račun potrošnje memorije. Na sreću, u slučaju 32-bitnih ključeva i vrijednosti, to se može opravdati. Ako u gornjem primjeru, u tabeli sa 128 miliona elemenata, zadržimo faktor iskorištenja od 0,25, onda u nju ne možemo postaviti više od 32 miliona elemenata, a preostalih 96 miliona slotova će biti izgubljeno - 8 bajtova za svaki par , 768 MB izgubljene memorije.
Napominjemo da govorimo o gubitku memorije video kartice, koja je vrijedniji resurs od sistemske memorije. Iako većina modernih desktop grafičkih kartica koje podržavaju CUDA-u imaju najmanje 4 GB memorije (u trenutku pisanja NVIDIA 2080 Ti ima 11 GB), ipak ne bi bila najmudrija odluka izgubiti takve količine.
Kasnije ću pisati više o kreiranju hash tablica za video kartice koje nemaju problema s dubinom sondiranja, kao i o načinima ponovnog korištenja izbrisanih slotova.
Merenje dubine sondiranja
Da bismo odredili dubinu ispitivanja ključa, možemo izdvojiti heš ključa (njegovog idealnog indeksa tabele) iz njegovog stvarnog indeksa tabele:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Zbog magije komplementarnih binarnih brojeva dva i činjenice da je kapacitet heš tabele dva na stepen dva, ovaj pristup će raditi čak i kada se indeks ključa pomeri na početak tabele. Uzmimo ključ koji je heširan na 1, ali je umetnut u slot 3. Tada za tabelu kapaciteta 4 dobijamo (3 — 1) & 3, što je ekvivalentno 2.
zaključak
Ako imate pitanja ili komentara, pošaljite mi e-poštu na ili otvorite novu temu.
Ovaj kod je napisan inspirisan odličnim člancima:
U budućnosti ću nastaviti pisati o implementaciji hash tablica za video kartice i analizirati njihove performanse. Moji planovi uključuju ulančavanje, Robin Hood heširanje i cuckoo heširanje korištenjem atomskih operacija u strukturama podataka koje su prilagođene GPU-u.
izvor: www.habr.com