Postitasin selle Githubisse .
See on lihtne GPU räsitabel, mis suudab töödelda sadu miljoneid sisestusi sekundis. Minu NVIDIA GTX 1060 sülearvutisse lisab kood 64 miljonit juhuslikult genereeritud võtmeväärtuste paari umbes 210 ms jooksul ja eemaldab 32 miljonit paari umbes 64 ms jooksul.
See tähendab, et sülearvuti kiirus on ligikaudu 300 miljonit lisamist sekundis ja 500 miljonit kustutamist sekundis.
Tabel on kirjutatud CUDA keeles, kuigi sama tehnikat saab rakendada ka HLSL-i või GLSL-i puhul. Rakendusel on videokaardi suure jõudluse tagamiseks mitmeid piiranguid:
- Töödeldakse ainult 32-bitiseid võtmeid ja samu väärtusi.
- Räsitabel on kindla suurusega.
- Ja see suurus peab olema võrdne kahe võimsusega.
Võtmete ja väärtuste jaoks peate reserveerima lihtsa eraldusmärgi (ülaltoodud koodis on see 0xffffffff).
Räsilaud ilma lukkudeta
Räsitabel kasutab avatud adresseerimist koos , see tähendab, et see on lihtsalt mällu salvestatud võtme-väärtuste paaride massiiv, millel on suurepärane vahemälu jõudlus. Sama ei saa öelda aheldamise kohta, mis hõlmab lingitud loendist osuti otsimist. Räsitabel on lihtne massiiv, mis salvestab elemente KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Tabeli suurus on kahe aste, mitte algarv, sest maski pow2/AND rakendamiseks piisab ühest kiirest käsust, kuid moodulite operaator on palju aeglasem. See on oluline lineaarse sondeerimise puhul, kuna lineaarses tabeliotsingus peab pesa indeks olema mähitud igasse pilusse. Selle tulemusena lisandub igas pesas mooduli maksumus operatsioonile.
Tabel salvestab ainult iga elemendi võtme ja väärtuse, mitte võtme räsi. Kuna tabelis on ainult 32-bitised võtmed, arvutatakse räsi väga kiiresti. Ülaltoodud kood kasutab Murmur3 räsi, mis teeb vaid mõne nihke, XOR-i ja korrutamise.
Räsitabel kasutab lukustuskaitse tehnikaid, mis ei sõltu mälujärjekorrast. Isegi kui mõned kirjutamistoimingud häirivad teiste selliste toimingute järjekorda, säilitab räsitabel ikkagi õige oleku. Sellest räägime allpool. Tehnika töötab suurepäraselt videokaartidega, mis jooksevad samaaegselt tuhandeid lõime.
Räsitabelis olevad võtmed ja väärtused lähtestatakse tühjaks.
Koodi saab muuta nii, et see käsitleks ka 64-bitisi võtmeid ja väärtusi. Võtmed nõuavad aatomi lugemist, kirjutamist ning võrdlemise ja vahetamise toiminguid. Ja väärtused nõuavad aatomi lugemis- ja kirjutamisoperatsioone. Õnneks on CUDA-s 32- ja 64-bitiste väärtuste lugemis-kirjutamistoimingud atomaarsed seni, kuni need on loomulikult joondatud (vt allpool). ) ja kaasaegsed videokaardid toetavad 64-bitiseid aatomite võrdlemise ja vahetamise toiminguid. Loomulikult väheneb jõudlus 64-bitisele üleminekul veidi.
Räsitabeli olek
Igal räsitabelis oleva võtme-väärtuse paaril võib olla üks neljast olekust:
- Võti ja väärtus on tühjad. Selles olekus lähtestatakse räsitabel.
- Võti on kirja pandud, aga väärtust pole veel kirja pandud. Kui mõni teine lõim loeb praegu andmeid, naaseb see tühjalt. See on normaalne, sama oleks juhtunud, kui mõni teine täitmislõng oleks veidi varem töötanud ja me räägime samaaegsest andmestruktuurist.
- Nii võti kui ka väärtus salvestatakse.
- Väärtus on saadaval teistele täitmislõimedele, kuid võti pole veel saadaval. See võib juhtuda, kuna CUDA programmeerimismudelil on lõdvalt järjestatud mälumudel. See on normaalne; igal juhul on võti endiselt tühi, isegi kui väärtus enam ei ole.
Oluline nüanss on see, et kui võti on pessa kirjutatud, siis see enam ei liigu – isegi kui võti kustutatakse, räägime sellest allpool.
Räsitabeli kood töötab isegi lõdvalt järjestatud mälumudelitega, milles mälu lugemise ja kirjutamise järjekord pole teada. Kui vaatame räsitabelis sisestamist, otsimist ja kustutamist, pidage meeles, et iga võtme-väärtuse paar on ühes neljast ülalkirjeldatud olekust.
Räsitabelisse sisestamine
CUDA funktsioon, mis lisab võtme-väärtuse paarid räsitabelisse, näeb välja järgmine:
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);
}
}
Võtme sisestamiseks kordab kood läbi räsitabeli massiivi, alustades sisestatud võtme räsist. Iga massiivi pesa teostab aatomi võrdlemise ja vahetamise toimingu, mis võrdleb selles pesas olevat võtit tühjaks. Kui tuvastatakse mittevastavus, värskendatakse pesas olevat võtit sisestatud võtmega ja seejärel tagastatakse algne pesa võti. Kui see algne võti oli tühi või vastas sisestatud võtmele, leidis kood sisestamiseks sobiva pesa ja sisestas sisestatud väärtuse pessa.
Kui ühes kerneli kõnes gpu_hashtable_insert() sama võtmega on mitu elementi, siis saab võtmepessa kirjutada nende mis tahes väärtuse. Seda peetakse normaalseks: kõne ajal üks võtmeväärtuse kirjutamine õnnestub, kuid kuna see kõik toimub paralleelselt mitmes täitmise lõimes, ei saa me ennustada, milline mällu kirjutamine on viimane.
Räsitabeli otsing
Kood võtmete otsimiseks:
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);
}
}
Tabelisse salvestatud võtme väärtuse leidmiseks kordame massiivi, alustades otsitava võtme räsist. Igas pesas kontrollime, kas võti on see, mida otsime, ja kui jah, siis tagastame selle väärtuse. Samuti kontrollime, kas võti on tühi, ja kui jah, siis katkestame otsingu.
Kui me võtit ei leia, tagastab kood tühja väärtuse.
Kõiki neid otsingutoiminguid saab sisestamise ja kustutamise kaudu teha samaaegselt. Igal tabeli paaril on üks neljast ülalkirjeldatud olekust voo jaoks.
Kustutamine räsitabelis
Võtmete kustutamise kood:
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);
}
}
Võtme kustutamine toimub ebatavalisel viisil: jätame võtme tabelisse ja märgime selle väärtuse (mitte võtme enda) tühjaks. See kood on väga sarnane lookup(), välja arvatud see, et kui võtmel leitakse vaste, muudab see selle väärtuse tühjaks.
Nagu eespool mainitud, kui võti on pesasse kirjutatud, siis seda enam ei liigutata. Isegi kui element tabelist kustutatakse, jääb võti paigale, selle väärtus muutub lihtsalt tühjaks. See tähendab, et me ei pea pesa väärtuse jaoks kasutama aatomikirjutusoperatsiooni, sest pole vahet, kas praegune väärtus on tühi või mitte – see jääb ikkagi tühjaks.
Räsitabeli suuruse muutmine
Räsitabeli suurust saab muuta, luues suurema tabeli ja sisestades sinna mittetühjad elemendid vanast tabelist. Ma ei rakendanud seda funktsiooni, kuna soovisin, et näidiskood oleks lihtne. Veelgi enam, CUDA programmides tehakse mälu eraldamine sageli pigem hostikoodis kui CUDA tuumas.
Artikkel kirjeldab, kuidas muuta sellist lukuga kaitstud andmestruktuuri.
Konkurentsivõime
Ülaltoodud funktsiooni koodijuppides gpu_hashtable_insert(), _lookup() и _delete() töödelda ühte võtme-väärtuse paari korraga. Ja madalam gpu_hashtable_insert(), _lookup() и _delete() töödelda paralleelselt paaride massiivi, iga paar eraldi GPU täitmislõimes:
// 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);
}
}
Lukustuskindel räsitabel toetab samaaegseid lisamisi, otsinguid ja kustutamisi. Kuna võtme-väärtuste paarid on alati ühes neljast olekust ja võtmed ei liigu, garanteerib tabel õigsuse ka siis, kui samaaegselt kasutatakse erinevat tüüpi toiminguid.
Kui aga töötleme sisestusi ja kustutamisi paralleelselt ja kui paaride sisendmassiivis on dubleerivad võtmed, siis ei saa me ennustada, millised paarid “võidavad” – kirjutatakse räsitabelisse viimasena. Oletame, et kutsusime sisestuskoodi paaride sisendmassiiviga A/0 B/1 A/2 C/3 A/4. Kui kood on valmis, paaritatakse B/1 и C/3 on tabelis garanteeritud, kuid samal ajal ilmub sinna ükskõik milline paar A/0, A/2 või A/4. See võib olla probleem, aga ei pruugi – kõik sõltub rakendusest. Võite ette teada, et sisendmassiivis ei ole dubleerivaid võtmeid või ei pruugi teid huvitada, milline väärtus kirjutati viimati.
Kui see on teie jaoks probleem, peate duplikaatpaarid eraldama erinevateks CUDA süsteemikõnedeks. CUDA-s lõpevad kõik kernelit kutsuvad toimingud alati enne järgmist kerneli kutset (vähemalt ühe lõime piires. Erinevates lõimedes käivitatakse tuumad paralleelselt). Ülaltoodud näites, kui helistate ühele kernelile rakendusega A/0 B/1 A/2 C/3, ja teine koos A/4, siis võti A saab väärtuse 4.
Nüüd räägime sellest, kas funktsioonid peaksid seda tegema lookup() и delete() kasutage tavalist või muutlikku kursorit räsitabelis paaride massiivile. teatab, et:
Kompilaator võib optimeerida lugemist ja kirjutamist globaalsesse või jagatud mällu... Need optimeerimised saab keelata märksõnaga
volatile: ... iga viide sellele muutujale kompileeritakse reaalseks mälu lugemis- või kirjutamiskäsku.
Õigsuse kaalutlused ei nõua rakendamist volatile. Kui täitmislõim kasutab vahemällu salvestatud väärtust varasemast lugemistoimingust, kasutab see veidi vananenud teavet. Kuid ikkagi on see teave räsitabeli õigest olekust teatud tuumakutse hetkel. Kui teil on vaja kasutada uusimat teavet, võite kasutada registrit volatile, kuid siis jõudlus veidi langeb: minu testide järgi langes 32 miljoni elemendi kustutamisel kiirus 500 miljonilt kustutamiselt/sek 450 miljonile kustutamisele/sek.
Производительность
64 miljoni elemendi sisestamise ja neist 32 miljoni kustutamise testis toimub konkurents std::unordered_map ja GPU jaoks praktiliselt puudub räsitabel:
std::unordered_map kulutas 70 691 ms elementide sisestamiseks ja eemaldamiseks ning seejärel vabastamiseks unordered_map (miljonitest elementidest vabanemine võtab palju aega, sest sees unordered_map tehakse mitu mälueraldust). Ausalt öeldes, std:unordered_map täiesti erinevad piirangud. See on üks protsessori täitmislõng, toetab mis tahes suurusega võtmeväärtusi, toimib hästi kõrge kasutusmäära juures ja näitab stabiilset jõudlust pärast mitut kustutamist.
GPU ja programmidevahelise suhtluse räsitabeli kestus oli 984 ms. See hõlmab aega, mis kulub tabeli mällu paigutamisele ja kustutamisele (ühekordne 1 GB mälu eraldamine, mis võtab CUDA-s veidi aega), elementide sisestamine ja kustutamine ning nende üle itereerimine. Arvesse lähevad ka kõik koopiad videokaardi mällu ja sealt.
Räsitabeli enda valmimine võttis aega 271 ms. See hõlmab aega, mis kulub videokaardil elementide sisestamiseks ja kustutamiseks, ega võta arvesse aega, mis kulub mällu kopeerimisele ja tulemuseks oleva tabeli itereerimisele. Kui GPU tabel elab pikka aega või kui räsitabel sisaldub täielikult videokaardi mälus (näiteks selleks, et luua räsitabel, mida kasutab muu GPU kood, mitte keskprotsessor), testi tulemus on asjakohane.
Videokaardi räsitabel näitab kõrget jõudlust tänu suurele läbilaskevõimele ja aktiivsele paralleelsusele.
Piirangud
Räsitabeli arhitektuuril on mõned probleemid, millest peaksite teadma:
- Lineaarset sondeerimist takistab rühmitamine, mis põhjustab tabelis olevate võtmete ebatäiuslikku paigutust.
- Funktsiooniga klahve ei eemaldata
deleteja aja jooksul ajavad nad laua segamini.
Selle tulemusel võib räsitabeli jõudlus järk-järgult halveneda, eriti kui see eksisteerib pikka aega ning sisaldab arvukalt lisamisi ja kustutamisi. Üks võimalus nende puuduste leevendamiseks on uuesti räsimine uude tabelisse, mille kasutusmäär on üsna madal ja eemaldatud võtmed uuesti räsimise käigus välja filtreerida.
Kirjeldatud probleemide illustreerimiseks kasutan ülaltoodud koodi, et luua 128 miljoni elemendiga tabel ja läbida 4 miljonit elementi, kuni olen täitnud 124 miljonit pesa (kasutusmäär umbes 0,96). Siin on tulemuste tabel, iga rida on CUDA kerneli kutse, et lisada ühte räsitabelisse 4 miljonit uut elementi:
Kasutusmäär
Sisestamise kestus 4 194 304 elementi
0,00
11,608448 ms (361,314798 miljonit klahvi sekundis)
0,03
11,751424 ms (356,918799 miljonit klahvi sekundis)
0,06
11,942592 ms (351,205515 miljonit klahvi sekundis)
0,09
12,081120 ms (347,178429 miljonit klahvi sekundis)
0,12
12,242560 ms (342,600233 miljonit klahvi sekundis)
0,16
12,396448 ms (338,347235 miljonit klahvi sekundis)
0,19
12,533024 ms (334,660176 miljonit klahvi sekundis)
0,22
12,703328 ms (330,173626 miljonit klahvi sekundis)
0,25
12,884512 ms (325,530693 miljonit klahvi sekundis)
0,28
13,033472 ms (321,810182 miljonit klahvi sekundis)
0,31
13,239296 ms (316,807174 miljonit klahvi sekundis)
0,34
13,392448 ms (313,184256 miljonit klahvi sekundis)
0,37
13,624000 ms (307,861434 miljonit klahvi sekundis)
0,41
13,875520 ms (302,280855 miljonit klahvi sekundis)
0,44
14,126528 ms (296,909756 miljonit klahvi sekundis)
0,47
14,399328 ms (291,284699 miljonit klahvi sekundis)
0,50
14,690304 ms (285,515123 miljonit klahvi sekundis)
0,53
15,039136 ms (278,892623 miljonit klahvi sekundis)
0,56
15,478656 ms (270,973402 miljonit klahvi sekundis)
0,59
15,985664 ms (262,379092 miljonit klahvi sekundis)
0,62
16,668673 ms (251,627968 miljonit klahvi sekundis)
0,66
17,587200 ms (238,486174 miljonit klahvi sekundis)
0,69
18,690048 ms (224,413765 miljonit klahvi sekundis)
0,72
20,278816 ms (206,831789 miljonit klahvi sekundis)
0,75
22,545408 ms (186,038058 miljonit klahvi sekundis)
0,78
26,053312 ms (160,989275 miljonit klahvi sekundis)
0,81
31,895008 ms (131,503463 miljonit klahvi sekundis)
0,84
42,103294 ms (99,619378 miljonit klahvi sekundis)
0,87
61,849056 ms (67,815164 miljonit klahvi sekundis)
0,90
105,695999 ms (39,682713 miljonit klahvi sekundis)
0,94
240,204636 ms (17,461378 miljonit klahvi sekundis)
Kasutamise suurenedes väheneb jõudlus. See pole enamikul juhtudel soovitav. Kui rakendus lisab tabelisse elemente ja seejärel loobub (näiteks raamatus sõnade lugemisel), siis pole see probleem. Kui aga rakendus kasutab pikaealist räsitabelit (näiteks graafikaredaktoris piltide mittetühjade osade salvestamiseks, kuhu kasutaja sageli teavet lisab ja kustutab), võib see käitumine olla problemaatiline.
Ja mõõtis räsitabeli proovimise sügavust pärast 64 miljonit lisamist (kasutustegur 0,5). Keskmine sügavus oli 0,4774, seega oli enamik klahve kas parimas võimalikus pesas või parimast asendist ühe pilu kaugusel. Maksimaalne sondeerimissügavus oli 60.
Seejärel mõõtsin sondeerimissügavust 124 miljoni sisetükiga laual (kasutustegur 0,97). Keskmine sügavus oli juba 10,1757 ja maksimaalne - 6474 (!!). Lineaarse anduri jõudlus langeb märkimisväärselt kõrge kasutusmäära korral.
Parim on hoida selle räsitabeli kasutusmäär madalal. Kuid siis suurendame jõudlust mälutarbimise arvelt. Õnneks on 32-bitiste võtmete ja väärtuste puhul see õigustatud. Kui ülaltoodud näites jätame 128 miljoni elemendiga tabelis kasutusteguriks 0,25, siis ei saa me sinna paigutada rohkem kui 32 miljonit elementi ja ülejäänud 96 miljonit pesa lähevad kaotsi - 8 baiti iga paari kohta , 768 MB mälu on kadunud.
Pange tähele, et me räägime videokaardi mälu kadumisest, mis on väärtuslikum ressurss kui süsteemimälu. Kuigi enamikul kaasaegsetel lauaarvuti graafikakaartidel, mis toetavad CUDAt, on vähemalt 4 GB mälu (kirjutamise hetkel on NVIDIA 2080 Ti-l 11 GB), poleks siiski kõige targem otsus selliseid summasid kaotada.
Hiljem kirjutan lähemalt räsitabelite loomisest videokaartidele, millel ei ole sondeerimissügavusega probleeme, samuti kustutatud pesade taaskasutamise võimalustest.
Helisügavuse mõõtmine
Võtme uurimissügavuse määramiseks saame eraldada võtme räsi (selle ideaalse tabeliindeksi) selle tegelikust tabeliindeksist:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Kahe kahe komplementaarsete kahendarvude võlu ja asjaolu tõttu, et räsitabeli maht on kaks kahe astmega, töötab see lähenemisviis isegi siis, kui võtmeindeks on viidud tabeli algusesse. Võtame võtme, mille räsi on 1, kuid mis on sisestatud pesasse 3. Seejärel saame 4 mahuga tabeli jaoks (3 — 1) & 3, mis võrdub 2-ga.
Järeldus
Kui teil on küsimusi või kommentaare, saatke mulle e-kiri aadressil või ava uus teema.
See kood on kirjutatud suurepärastest artiklitest inspireerituna:
Edaspidi jätkan videokaartide räsitabelite juurutuste kirjutamist ja nende jõudluse analüüsimist. Minu plaanid hõlmavad aheldamist, Robin Hoodi räsimist ja kägu räsimist, kasutades aatomoperatsioone GPU-sõbralikes andmestruktuurides.
Allikas: www.habr.com