Objavil sem ga na Githubu .
To je preprosta zgoščena tabela GPU, ki lahko obdela stotine milijonov vstavkov na sekundo. Na mojem prenosnem računalniku NVIDIA GTX 1060 koda v približno 64 ms vstavi 210 milijonov naključno ustvarjenih parov ključ-vrednost in odstrani 32 milijonov parov v približno 64 ms.
To pomeni, da je hitrost na prenosnem računalniku približno 300 milijonov vstavljanj/s in 500 milijonov brisanja/s.
Tabela je napisana v CUDA, čeprav se lahko ista tehnika uporabi za HLSL ali GLSL. Izvedba ima več omejitev za zagotavljanje visoke zmogljivosti na video kartici:
- Obdelujejo se samo 32-bitni ključi in enake vrednosti.
- Zgoščevalna tabela ima fiksno velikost.
- In ta velikost mora biti enaka dve na potenco.
Za ključe in vrednosti morate rezervirati preprosto ločilno oznako (v zgornji kodi je to 0xffffffff).
Hash tabela brez ključavnic
Zgoščena tabela uporablja odprto naslavljanje z , to je preprosto niz parov ključ-vrednost, ki je shranjen v pomnilniku in ima vrhunsko zmogljivost predpomnilnika. Tega ne moremo reči za veriženje, ki vključuje iskanje kazalca na povezanem seznamu. Zgoščevalna tabela je preprosta matrika, ki shranjuje elemente KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Velikost tabele je potenca dvojke, ne praštevilo, ker je eno hitro navodilo dovolj za uporabo maske pow2/AND, vendar je operater modula veliko počasnejši. To je pomembno v primeru linearnega tipanja, saj mora biti pri iskanju linearne tabele indeks reže zavit v vsako režo. In kot rezultat, se strošek operacije doda modulo v vsako režo.
Tabela shrani samo ključ in vrednost za vsak element, ne pa zgoščene vrednosti ključa. Ker tabela hrani samo 32-bitne ključe, se zgoščevanje izračuna zelo hitro. Zgornja koda uporablja zgoščevanje Murmur3, ki izvede le nekaj premikov, XOR in množenja.
Zgoščevalna tabela uporablja zaščitne tehnike zaklepanja, ki so neodvisne od pomnilniškega vrstnega reda. Tudi če nekatere pisalne operacije zmotijo vrstni red drugih takih operacij, bo zgoščena tabela še vedno ohranila pravilno stanje. O tem bomo govorili spodaj. Tehnika odlično deluje z grafičnimi karticami, ki hkrati izvajajo na tisoče niti.
Ključi in vrednosti v zgoščeni tabeli so inicializirani tako, da so prazni.
Kodo je mogoče spremeniti tako, da obravnava tudi 64-bitne ključe in vrednosti. Ključi zahtevajo atomske operacije branja, pisanja in primerjave in zamenjave. In vrednosti zahtevajo atomske operacije branja in pisanja. Na srečo so v CUDA operacije branja in pisanja za 32- in 64-bitne vrednosti atomične, dokler so naravno poravnane (glejte spodaj). ), sodobne video kartice pa podpirajo 64-bitne atomske operacije primerjave in izmenjave. Seveda se bo pri prehodu na 64 bitov zmogljivost nekoliko zmanjšala.
Stanje zgoščene tabele
Vsak par ključ-vrednost v zgoščeni tabeli ima lahko eno od štirih stanj:
- Ključ in vrednost sta prazna. V tem stanju je zgoščena tabela inicializirana.
- Ključ je zapisan, vrednost pa še ni zapisana. Če druga nit trenutno bere podatke, se vrne prazna. To je normalno, enako bi se zgodilo, če bi druga nit izvajanja delovala malo prej in govorimo o sočasni strukturi podatkov.
- Zabeležita se tako ključ kot vrednost.
- Vrednost je na voljo drugim nitim izvajanja, vendar ključ še ni. To se lahko zgodi, ker ima programski model CUDA ohlapno urejen spominski model. To je normalno; v vsakem primeru je ključ še vedno prazen, tudi če vrednost ni več to.
Pomemben odtenek je, da ko je ključ zapisan v režo, se ne premakne več - tudi če je ključ izbrisan, bomo o tem govorili spodaj.
Koda razpršilne tabele deluje celo z ohlapno urejenimi pomnilniškimi modeli, v katerih vrstni red branja in zapisovanja pomnilnika ni znan. Ko gledamo vstavljanje, iskanje in brisanje v zgoščevalni tabeli, ne pozabite, da je vsak par ključ-vrednost v enem od štirih zgoraj opisanih stanj.
Vstavljanje v razpršilno tabelo
Funkcija CUDA, ki vstavi pare ključ-vrednost v zgoščeno tabelo, je videti takole:
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);
}
}
Za vstavljanje ključa se koda ponovi skozi matriko zgoščene tabele, začenši z zgoščeno vrednostjo vstavljenega ključa. Vsaka reža v matriki izvede operacijo atomske primerjave in zamenjave, ki primerja ključ v tej reži s praznim. Če se odkrije neujemanje, se ključ v reži posodobi z vstavljenim ključem, nato pa se vrne originalni ključ reže. Če je bil ta prvotni ključ prazen ali se je ujemal z vstavljenim ključem, je koda našla primerno režo za vstavljanje in v režo vstavila vstavljeno vrednost.
Če v enem klicu jedra gpu_hashtable_insert() obstaja več elementov z istim ključem, potem lahko katero koli njihovo vrednost zapišete v režo za ključ. To velja za normalno: eden od zapisov ključ-vrednost med klicem bo uspel, a ker se vse to dogaja vzporedno v več nitih izvajanja, ne moremo predvideti, kateri zapis v pomnilnik bo zadnji.
Iskanje zgoščene tabele
Koda za iskanje ključev:
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 bi našli vrednost ključa, shranjenega v tabeli, ponovimo matriko, začenši z zgoščeno vrednostjo ključa, ki ga iščemo. V vsaki reži preverimo, ali je ključ tisti, ki ga iščemo, in če je, vrnemo njegovo vrednost. Preverimo tudi, ali je ključ prazen, in če je, prekinemo iskanje.
Če ključa ne najdemo, koda vrne prazno vrednost.
Vse te iskalne operacije je mogoče izvajati sočasno z vstavljanjem in brisanjem. Vsak par v tabeli bo imel eno od štirih zgoraj opisanih stanj za tok.
Brisanje v zgoščeni tabeli
Koda za brisanje ključev:
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 poteka na nenavaden način: ključ pustimo v tabeli in njegovo vrednost (ne ključa samega) označimo kot prazno. Ta koda je zelo podobna lookup(), le da ko je na ključu najdeno ujemanje, postane njegova vrednost prazna.
Kot je navedeno zgoraj, ko je ključ zapisan v režo, se ne premakne več. Tudi ko je element izbrisan iz tabele, ključ ostane na mestu, njegova vrednost preprosto postane prazna. To pomeni, da nam ni treba uporabiti operacije atomskega pisanja za vrednost reže, ker ni pomembno, ali je trenutna vrednost prazna ali ne – še vedno bo prazna.
Spreminjanje velikosti zgoščene tabele
Velikost zgoščene tabele lahko spremenite tako, da ustvarite večjo tabelo in vanjo vstavite neprazne elemente iz stare tabele. Te funkcije nisem implementiral, ker sem želel ohraniti vzorčno kodo preprosto. Poleg tega se v programih CUDA dodelitev pomnilnika pogosto izvede v gostiteljski kodi in ne v jedru CUDA.
V članku opisuje, kako spremeniti tako z zaklepanjem zaščiteno podatkovno strukturo.
Tekmovalnost
V zgornjih delčkih kode funkcije gpu_hashtable_insert(), _lookup() и _delete() obdelujejo en par ključ-vrednost naenkrat. In nižje gpu_hashtable_insert(), _lookup() и _delete() vzporedno obdelajte niz parov, vsak par v ločeni izvršilni niti GPE:
// 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);
}
}
Zgoščevalna tabela, odporna na zaklepanje, podpira sočasne vstavke, iskanja in brisanja. Ker so pari ključ-vrednost vedno v enem od štirih stanj in se ključi ne premikajo, tabela zagotavlja pravilnost, tudi če so različne vrste operacij uporabljene hkrati.
Če pa vzporedno obdelamo serijo vstavljanj in izbrisov in če vhodna matrika parov vsebuje podvojene ključe, potem ne bomo mogli predvideti, kateri pari bodo "zmagali" - zadnji bodo zapisani v zgoščevalno tabelo. Recimo, da smo kodo za vstavljanje poklicali z vhodno matriko parov A/0 B/1 A/2 C/3 A/4. Ko je koda končana, pari B/1 и C/3 so zagotovljeno prisotni v tabeli, hkrati pa se bo v njej pojavil kateri koli od parov A/0, A/2 ali A/4. To je lahko težava ali pa tudi ne – vse je odvisno od aplikacije. Morda že vnaprej veste, da v vhodni matriki ni podvojenih ključev, ali pa vam je vseeno, katera vrednost je bila zapisana zadnja.
Če je to za vas težava, morate podvojene pare ločiti v različne sistemske klice CUDA. V CUDA se vsaka operacija, ki kliče jedro, vedno zaključi pred naslednjim klicem jedra (vsaj znotraj ene niti. V različnih nitih se jedra izvajajo vzporedno). V zgornjem primeru, če pokličete eno jedro z A/0 B/1 A/2 C/3, drugi pa z A/4, nato ključ A bo dobil vrednost 4.
Zdaj pa se pogovorimo o tem, ali bi morale funkcije lookup() и delete() uporabite navaden ali nestanoviten kazalec na niz parov v zgoščevalni tabeli. navaja, da:
Prevajalnik se lahko odloči za optimizacijo branja in pisanja v globalni ali skupni pomnilnik... Te optimizacije lahko onemogočite s ključno besedo
volatile: ... vsako sklicevanje na to spremenljivko se prevede v navodilo za branje ali pisanje pravega pomnilnika.
Premisleki o pravilnosti ne zahtevajo uporabe volatile. Če izvajalna nit uporablja predpomnjeno vrednost iz prejšnje operacije branja, bo uporabljala rahlo zastarele informacije. Še vedno pa je to informacija iz pravilnega stanja zgoščene tabele v določenem trenutku klica jedra. Če želite uporabiti najnovejše informacije, lahko uporabite kazalo volatile, potem pa se bo zmogljivost nekoliko zmanjšala: po mojih testih se je pri brisanju 32 milijonov elementov hitrost zmanjšala s 500 milijonov izbrisov/s na 450 milijonov izbrisov/s.
Produktivnost
V testu vstavljanja 64 milijonov elementov in brisanja 32 milijonov jih je tekmovalnost med std::unordered_map in praktično ni zgoščevalne tabele za GPE:
std::unordered_map porabil 70 ms za vstavljanje in odstranjevanje elementov ter njihovo sprostitev unordered_map (da se znebite milijonov elementov vzame veliko časa, ker znotraj unordered_map izvede se več dodelitev pomnilnika). Iskreno povedano, std:unordered_map popolnoma drugačne omejitve. To je nit izvajanja z enim procesorjem, podpira ključne vrednosti poljubne velikosti, dobro deluje pri visokih stopnjah izkoriščenosti in kaže stabilno delovanje po večkratnem izbrisu.
Trajanje zgoščene tabele za GPE in medprogramsko komunikacijo je bilo 984 ms. To vključuje čas, porabljen za postavitev tabele v pomnilnik in njeno brisanje (enkratna dodelitev 1 GB pomnilnika, kar traja nekaj časa v CUDA), vstavljanje in brisanje elementov ter ponavljanje po njih. Upoštevane so tudi vse kopije v in iz pomnilnika video kartice.
Sama zgoščena tabela je dokončala 271 ms. To vključuje čas, ki ga video kartica porabi za vstavljanje in brisanje elementov, in ne upošteva časa, porabljenega za kopiranje v pomnilnik in ponavljanje po dobljeni tabeli. Če tabela GPU obstaja dolgo časa ali če je tabela zgoščenih vrednosti v celoti shranjena v pomnilniku video kartice (na primer za ustvarjanje tabele zgoščenih vrednosti, ki jo bo uporabljala druga koda GPE in ne osrednji procesor), potem rezultat testa je pomemben.
Zgoščevalna tabela za video kartico kaže visoko zmogljivost zaradi visoke prepustnosti in aktivne paralelizacije.
Omejitve
Arhitektura zgoščevalne tabele ima nekaj težav, na katere se morate zavedati:
- Linearno tipanje ovira združevanje v gruče, zaradi česar so ključi v tabeli postavljeni manj kot popolno.
- S funkcijo se tipke ne odstranijo
deletein sčasoma zamašijo mizo.
Posledično se lahko učinkovitost zgoščevalne tabele postopoma poslabša, še posebej, če obstaja dlje časa in ima številne vstavke in brisanja. Eden od načinov za ublažitev teh pomanjkljivosti je ponovno zgoščevanje v novo tabelo z dokaj nizko stopnjo izkoriščenosti in filtriranje odstranjenih ključev med ponovnim zgoščevanjem.
Za ponazoritev opisanih težav bom uporabil zgornjo kodo, da ustvarim tabelo s 128 milijoni elementov in se pomikam po 4 milijonih elementov, dokler ne zapolnim 124 milijonov rež (stopnja izkoriščenosti približno 0,96). Tukaj je tabela rezultatov, vsaka vrstica je klic jedra CUDA za vstavljanje 4 milijonov novih elementov v eno zgoščeno tabelo:
Stopnja uporabe
Trajanje vstavljanja 4 elementov
0,00
11,608448 ms (361,314798 milijonov ključev/sek.)
0,03
11,751424 ms (356,918799 milijonov ključev/sek.)
0,06
11,942592 ms (351,205515 milijonov ključev/sek.)
0,09
12,081120 ms (347,178429 milijonov ključev/sek.)
0,12
12,242560 ms (342,600233 milijonov ključev/sek.)
0,16
12,396448 ms (338,347235 milijonov ključev/sek.)
0,19
12,533024 ms (334,660176 milijonov ključev/sek.)
0,22
12,703328 ms (330,173626 milijonov ključev/sek.)
0,25
12,884512 ms (325,530693 milijonov ključev/sek.)
0,28
13,033472 ms (321,810182 milijonov ključev/sek.)
0,31
13,239296 ms (316,807174 milijonov ključev/sek.)
0,34
13,392448 ms (313,184256 milijonov ključev/sek.)
0,37
13,624000 ms (307,861434 milijonov ključev/sek.)
0,41
13,875520 ms (302,280855 milijonov ključev/sek.)
0,44
14,126528 ms (296,909756 milijonov ključev/sek.)
0,47
14,399328 ms (291,284699 milijonov ključev/sek.)
0,50
14,690304 ms (285,515123 milijonov ključev/sek.)
0,53
15,039136 ms (278,892623 milijonov ključev/sek.)
0,56
15,478656 ms (270,973402 milijonov ključev/sek.)
0,59
15,985664 ms (262,379092 milijonov ključev/sek.)
0,62
16,668673 ms (251,627968 milijonov ključev/sek.)
0,66
17,587200 ms (238,486174 milijonov ključev/sek.)
0,69
18,690048 ms (224,413765 milijonov ključev/sek.)
0,72
20,278816 ms (206,831789 milijonov ključev/sek.)
0,75
22,545408 ms (186,038058 milijonov ključev/sek.)
0,78
26,053312 ms (160,989275 milijonov ključev/sek.)
0,81
31,895008 ms (131,503463 milijonov ključev/sek.)
0,84
42,103294 ms (99,619378 milijonov ključev/sek.)
0,87
61,849056 ms (67,815164 milijonov ključev/sek.)
0,90
105,695999 ms (39,682713 milijonov ključev/sek.)
0,94
240,204636 ms (17,461378 milijonov ključev/sek.)
Ko se izkoriščenost poveča, se zmogljivost zmanjša. To v večini primerov ni zaželeno. Če aplikacija vstavi elemente v tabelo in jih nato zavrže (na primer pri štetju besed v knjigi), potem to ni problem. Če pa aplikacija uporablja dolgotrajno zgoščeno tabelo (na primer v grafičnem urejevalniku za shranjevanje nepraznih delov slik, kjer uporabnik pogosto vnaša in briše informacije), potem je to vedenje lahko problematično.
In izmeril globino sondiranja razpršilne tabele po 64 milijonih vstavkov (faktor izkoriščenosti 0,5). Povprečna globina je bila 0,4774, tako da je bila večina ključev v najboljši možni reži ali eno režo stran od najboljšega položaja. Največja globina sondiranja je bila 60.
Globino tipanja sem nato izmeril na mizi s 124 milijoni ploščic (faktor izkoristka 0,97). Povprečna globina je bila že 10,1757, največja pa - 6474 (!!). Zmogljivost linearnega zaznavanja znatno pade pri visokih stopnjah uporabe.
Najbolje je, da je stopnja uporabe te zgoščene tabele nizka. Ampak potem povečamo zmogljivost na račun porabe pomnilnika. Na srečo je v primeru 32-bitnih ključev in vrednosti to mogoče upravičiti. Če v zgornjem primeru v tabeli s 128 milijoni elementov ohranimo faktor izkoriščenosti 0,25, potem vanjo ne smemo postaviti več kot 32 milijonov elementov, preostalih 96 milijonov slotov pa bo izgubljenih - 8 bajtov za vsak par , 768 MB izgubljenega pomnilnika.
Upoštevajte, da govorimo o izgubi pomnilnika video kartice, ki je dragocenejši vir kot sistemski pomnilnik. Čeprav ima večina sodobnih namiznih grafičnih kartic, ki podpirajo CUDA, vsaj 4 GB pomnilnika (v času pisanja ima NVIDIA 2080 Ti 11 GB), vseeno ne bi bila najbolj modra odločitev izgubiti takšne količine.
Kasneje bom napisal več o ustvarjanju zgoščevalnih tabel za grafične kartice, ki nimajo težav z globino sondiranja, pa tudi o načinih ponovne uporabe izbrisanih rež.
Merjenje globine sondiranja
Za določitev globine sondiranja ključa lahko izvlečemo zgoščeno vrednost ključa (njegov idealni indeks tabele) iz dejanskega indeksa tabele:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Zaradi čarovnije binarnih števil komplementa dveh in dejstva, da je zmogljivost zgoščevalne tabele enaka dva na potenco dva, bo ta pristop deloval, tudi če je ključni indeks premaknjen na začetek tabele. Vzemimo ključ, ki je zgoščen na 1, vendar je vstavljen v režo 3. Nato za tabelo s kapaciteto 4 dobimo (3 — 1) & 3, kar je enakovredno 2.
Zaključek
Če imate vprašanja ali komentarje, mi pošljite e-pošto na ali pa odpri novo temo v.
Ta koda je bila napisana po navdihu odličnih člankov:
V prihodnosti bom še naprej pisal o implementacijah zgoščenih tabel za grafične kartice in analiziral njihovo delovanje. Moji načrti vključujejo veriženje, Robin Hood zgoščevanje in kukavičasto zgoščevanje z uporabo atomskih operacij v podatkovnih strukturah, ki so prijazne GPU.
Vir: www.habr.com