Olen lukenud Githubiin uutta projektia A Simple GPU Hash Table.

Tämä on yksinkertainen GPU-hajautustaulu, joka pystyy käsittelemään satoja miljoonia lisäyksiä sekunnissa. Omalla kannettavalla tietokoneellani, jossa on NVIDIA GTX 1060, koodi lisää noin 64 miljoonaa satunnaisesti luotua avain-arvo-paria noin 210 ms:ssa ja poistaa 32 miljoonaa paria noin 64 ms:ssa.

Eli kannettavalla tietokoneella nopeus on noin 300 miljonaa lisäystä/sek ja 500 miljonaa poistoa/sek.

Taulukko on kirjoitettu CUDAlla, vaikka samaa menetelmää voidaan soveltaa HLSL:ään tai GLSL:ään. Toteutuksessa on useita rajoituksia, jotka varmistavat korkean suorituskyvyn näytönohjaimella:

• Käsitellään vain 32-bittisiä avaimia ja samankaltaisia arvoja.
• Hajautustaululla on kiinteä koko.
• Ja tämän koon tulisi olla kahden potenssi.

Avaimille ja arvoille on varattava yksinkertainen erottava merkki (esimerkkikoodissa tämä on 0xffffffff).

Ilman lukitusta oleva hajautustaulu

Hajautustaulussa käytetään avointa osoitusta lineaarisella tutkimisella, eli se on vain avain-arvo-parien taulukko, joka säilytetään muistissa ja jolla on erinomainen välimuistisuorituskyky. Tämä ei päde ketjutasoitukseen (chaining), joka tarkoittaa osoittimen etsimistä sidotussa listassa. Hajautustaulu on yksinkertainen taulukko, joka tallentaa elementit KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Taulukon koko on kahden potenssi, ei alkuluku, koska pow2/AND-maskeja voidaan käyttää yhdellä nopealla komennolla, kun taas moduulioperaattori toimii paljon hitaammin. Tämä on tärkeää lineaarisen tutkimisen yhteydessä, koska lineaarisessa haussa taulukon slotin indeksi on oltava kierretty jokaisessa slotissa. Tämän seurauksena moduulioperaation kustannus lisätään jokaiseen slotin.

Taulukko tallentaa vain avaimen ja arvon jokaiselle elementille, ei avaimen hajautusta. Koska taulukko tallentaa vain 32-bittisiä avaimia, hajautus lasketaan erittäin nopeasti. Esimerkkikoodissa käytetään Murmur3-hajautusta, joka suorittaa vain muutamia siirtoja, XOR:ia ja kertolaskuja.

Häkk-tabelis kasutatakse blokeerimise kaitse meetodit, mis ei sõltu mälu paigutuse järjekorrast. Isegi kui mõned kirjutusoperatsioonid rikuvad teiste selliste operatsioonide järjekorda, säilitab häkk-tabel siiski õige oleku. Sellest räägime allpool. Meetod töötab suurepäraselt graafikakaartidega, kus tuhanded voolud töötavad samal ajal.

Häkk-tabelis initsialiseeritakse võtmed ja väärtused tühjaks.

Koodi saab modifitseerida, et see saaks töödelda ka 64-bitiseid võtmeid ja väärtusi. Võtmete jaoks on vajalikud aatomilised lugemise, kirjutamise ja võrgu operatsioonid (compare-and-swap). Väärtuste jaoks on vajalikud aatomilised lugemise ja kirjutamise operatsioonid. Õnneks on CUDA-s 32- ja 64-bitiste väärtuste lugemis-kirjutamisoperatsioonid aatomilised, kuni need on õigesti joondatud (vt. siin), ja kaasaegsed graafikakaardid toetavad 64-bitiseid aatomilisi võrgu operatsioone. Loomulikult väheneb jõudlus 64-bitisele üleminekul veidi.

Häkk-tabeli olek

Igal võtme-väärtuse paaril häkk-tabelis võib olla üks neljast olekust:

• Võti ja väärtus on tühjad. Sellises olekus initsialiseeritakse häkk-tabel.
• Võti on kirjutatud, kuid väärtus ei ole veel. Kui sel ajal loeb andmeid teine täitmisvoog, tagastab see tühja väärtuse. See on normaalne, sama oleks juhtunud, kui teine täitmisvoog oleks töötanud natuke varem, ja räägime konkurentsivõimelisest andmestruktuurist.
• Nii võti kui ka väärtus on kirjutatud.
• Väärtus on teistele täitmisvoogudele saadaval, kuid võti ei ole veel. See võib juhtuda, kuna CUDA programmeerimisreduksioon eeldab nõrgalt korraldatud mälumudelit. See on normaalne, ühelgi juhul on võti endiselt tühi, isegi kui väärtus ei ole.

Oluline nüanss on see, et alates hetkest, mil võti kirjutatakse slotile, ei liigu see enam — isegi kui võti kustutatakse, räägime sellest allpool.

Häkk-tabeli kood töötab isegi nõrgalt korraldatud mälumudelites, kus lugemise ja kirjutamise järjekord mälu sees ei ole teada. Kui me arutame häkk-tabli lisamist, otsimist ja kustutamist, pidage meeles, et iga võtme-väärtuse paar on ühes neljast eespool kirjeldatud olekust.

Häkk-tabelisse lisamine

CUDA-funktsioon, mis lisab hargnitud tabelisse võtme-väärtuse paare, näeb välja selline:

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 iteratiivne kood hargnitud tabeli massiivi, alustades sisestatava võtme hash'ist. Igas massiivi slotis teostatakse atomaarne võrdlemise ja vahetamise operatsioon, kus antud sloti võtme võrdletakse tühjaga. Kui leitakse mittesobivus, uuendatakse slotis olev võtme sisestatava võtmega ning seejärel tagastatakse algne sloti võti. Kui see algne võti oli tühi või vastas sisestatavale võtmele, leidis kood sobiva sloti ja sisestab sloti vastava väärtuse.

Kui ühes tuumas kutsumises gpu_hashtable_insert() on mitu elementi sama võtmega, siis võib ükskõik milline nende väärtustest olla salvestatud võtme sloti. See on normaalne: üks võtme-väärtuse salvestamise operatsioon tuumakutsumise käigus õnnestub, kuid kuna see toimub paralleelselt mitme täitmisvoo raames, ei saa me ennustada, milline mälu salvestamise operatsioon on viimane.

Otsingu hargnitud tabelis

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);
        }
}

Võtme leidmiseks, mis on tabelis, iteratiivne kood massiivi, alustades otsitava võtme hash'ist. Igas slotis kontrollime, kas võtme on see, mida me otsime, ja kui jah, siis tagastame selle väärtuse. Samuti kontrollime, kas võti on tühi ja kui jah, katkestame otsingu.

Kui me ei suuda võtme leida, tagastab kood tühja väärtuse.

Kõiki neid otsingu operatsioone saab teostada üheaegselt sisestuste ja kustutuste käigus. Igal tabeli paaril on voos üks neljast ülaltoodud olekust.

Kustutamine hargnitud tabelis

Kood võtmete kustutamiseks:

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);
    }
}

Klahvi kustutamine toimub ebatavaliselt: jätame võtme tabelisse ja märgistame selle väärtuse (mitte võtme endi) tühjaks. See kood sarnaneb väga lookup(), välja arvatud, et võtme vastavuse tuvastamisel teeb see selle väärtuse tühjaks.

Nagu eespool mainitud, kui võtme slotti kirjutatakse, ei liigu see enam. Isegi kui tabelist element kustutatakse, jääb võtme koht muutumatuks, lihtsalt selle väärtus muutub tühjaks. See tähendab, et me ei pea kasutama slotiväärtuse kirjutamiseks aatomilist operatsiooni, kuna pole vahet, kas praegune väärtus on tühi või mitte — see muutub ikkagi tühjaks.

Hästorinti suurendamine

Hästorinti saab suurendada, luues suurema tabeli ja sisestades sellesse mitte-tühjad elemendid vanast tabelist. Ma ei rakendanud seda funktsionaalsust, kuna tahtsin säilitada koodi lihtsana. Veelgi enam, CUDA programmides toimub mälu eraldamine sageli host-koodis, mitte CUDA ydas.

Artiklis Bloqueeringu vaba ootevaba hash-tabel on kirjeldatud, kuidas sellist blokeerimise vastu kaitstud andmestruktuuri muuta.

Konkurents

Ülaltoodud koodilõikudes funktsioonid gpu_hashtable_insert(), _lookup() ja _delete() töötlevad ühe võtme-väärtuse paari korraga. Allpool gpu_hashtable_insert(), _lookup() ja _delete() töötlevad paaride massiivi paralleelselt, iga paari eraldi GPU töövoos:

// 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);
    }
}

Bloqueeringu vastu kaitstud hash-tabel toetab konkurentsivõimelisi sisestusi, otsinguid ja kustutamisi. Kuna võtme-väärtuse paarid on alati ühes neljast olekust ja võtmed ei liigu, tagab tabel õigsuse isegi erinevate operatsioonide samaaegsel kasutamisel.

Kuid kui töötleme paralleelselt sisestuste ja kustutuste paketti ning kui sisendmassiivis on korduvad võtmed, ei suuda me ennustada, millised paarid "võidavad" — mis kirjutatakse hash-tabelisse viimaseid. Oletame, et kutsusime sisestuskoodi, mille sisendmassiiviks on paarid A/0 B/1 A/2 C/3 A/4. Kui kood on lõpule viidud, on paarid B/1 ja C/3 garanteeritult tabelis, kuid samas võivad seal olla ka mistahes paarid A/0, A/2 või A/4See, see, see. This could be a problem or it may not—everything depends on the application. You might already know there's no duplicate keys in the input array, or it might not matter which value was recorded last.

If this is a problem for you, then you need to separate the duplicate pairs into different system CUDA calls. In CUDA, any operation with a kernel call always completes before the next kernel call (at least within a single thread. In different threads, kernels execute in parallel). If in the example above, one kernel is called with A/0 B/1 A/2 C/3, and another with A/4, then the key A will receive the value 4.

Now let's discuss whether the functions lookup() ja delete() should use a plain or volatile pointer to the array of pairs in the hash table. The CUDA documentation states that:

The compiler may optimize read and write operations to global or shared memory at its discretion… These optimizations can be disabled using the keyword volatile: … any reference to this variable compiles to an actual read or write instruction in memory.

Correctness considerations do not require application of volatile. If the execution thread uses a cached value from an earlier read operation, it means it will use slightly outdated information. However, it is still information from a valid state of the hash table at a specific kernel call moment. If you need to use the latest information, you can use a pointer volatile, but then performance may slightly decrease: according to my tests—when deleting 32 million elements, the speed dropped from 500 million deletes/sec to 450 million deletes/sec.

Tõhusus

In the test for inserting 64 million elements and deleting 32 million of them, competition between std::unordered_map and the hash table for the GPU is virtually nonexistent:

std::unordered_map it took 70,691 ms to insert and delete elements followed by the freeing up of unordered_map (freeing up millions of elements takes quite some time because numerous memory allocations happen inside). Honestly, regarding unordered_map std:unordered_map std:unordered_map täiesti teised piirangud. See on ainus CPU-tegevusvoog, mis toetab võtme-väärtuse paare mistahes suuruses, töötab hästi kõrgete kasutuskoefitsientide korral ning näitab pärast rohkete kustutamiste järel stabiilset jõudlust.

Hash-tabeli tööaeg GPU ja protsessidevahelistes suhtlustes oli 984 ms. See hõlmab aega, mis kulus tabeli mälu paigutamiseks ja selle kustutamiseks (ühekordne 1 GB mälu eraldamine, mis võtab CUDA-s aega), elementide sisestamine ja kustutamine, samuti nende vaheliteriseerimine. Samuti on arvesse võetud kõik koopiad mälust ja graafikakaardi mälu.

Hash-tabeli töötamine võttis 271 ms. See hõlmab aega, mille graafikakaart kulutab elementide sisestamiseks ja kustutamiseks, ja ei arvestata aega, mis kulub mälu kopeerimiseks ja saadud tabeli vaheliteriseerimiseks. Kui GPU-tabel elab kaua, või kui hash-tabel asub täielikult graafikakaardi mälus (näiteks hash-tabeli loomiseks, mida kasutatakse teise GPU-koodi jaoks, mitte keskpäeva protsessori jaoks), on testi tulemus asjakohane.

Hash-tabel graafikakaardi jaoks näitab kõrget jõudlust, tänu suurele läbilaskvusele ja aktiivsele paralleliseerimisele.

Puudused

Hash-tabeli arhitektuuril on mitmeid probleeme, mida tasub meeles pidada:

• Lineaarne sondimine on häiritud klasterdamise tõttu, mille tõttu paigutatakse võtmed tabelisse kaugele mitte ideaalselt.
• Võtmed ei kustutata funktsiooniga delete ja aja jooksul ummistavad tabelit.

Kuna selle tulemusena võib hash-tabeli jõudlus järk-järgult väheneda, eriti kui see eksisteerib kaua ja selles viiakse läbi rohkelt sisestusi ja kustutamisi. Üks viis nende puuduste leevendamiseks on uut tabelit alamhõivega, piisavalt madala kasutuskoefitsiendiga ja kustutatud võtmete filtreerimisega alamhõive ajal.

Ülaltoodud probleemide illustreerimiseks kasutan eespool esitatud koodi, et luua 128 miljoni elemendi tabel, sisestan tsükliliselt 4 miljonit elementi, kuni olen täitnud 124 miljoni pesa (kasutuskoefitsient umbes 0,96). Siin on tulemuste tabel, iga rida on CUDA tuuma kutse 4 miljoni uue elemendi sisestamiseks ühte hash-tabelisse:

Kasutuskoefitsient
Sisestamisaeg 4 194 304 elementi jaoks

0,00
11,608448 ms (361,314798 miljonit võtit/s)

0,03
11,751424 ms (356,918799 miljonit võtit/s)

0,06
11,942592 ms (351,205515 miljonit võtit/s)

0,09
12,081120 ms (347,178429 miljonit võtit/s)

0,12
12,242560 ms (342,600233 miljonit võtit/s)

0,16
12,396448 ms (338,347235 miljonit võtit/s)

0,19
12,533024 ms (334,660176 miljonit võtit/s)

0,22
12,703328 ms (330,173626 miljonit võtit/s)

0,25
12,884512 ms (325,530693 miljonit võtit/s)

0,28
13,033472 ms (321,810182 miljonit võtit/s)

0,31
13,239296 ms (316,807174 miljonit võtit/s)

0,34
13,392448 ms (313,184256 miljonit võtit/s)

0,37
13,624000 ms (307,861434 miljonit võtit/s)

0,41
13,875520 ms (302,280855 miljonit võtit/s)

0,44
14,126528 ms (296,909756 miljonit võtit/s)

0,47
14,399328 ms (291,284699 miljonit võtit/s)

0,50
14,690304 ms (285,515123 miljonit võtit/s)

0,53
15,039136 ms (278,892623 miljonit võtit/s)

0,56
15,478656 ms (270,973402 miljonit võtit/s)

0,59
15,985664 ms (262,379092 miljonit võtit/s)

0,62
16,668673 ms (251,627968 miljonit võtit/s)

0,66
17,587200 ms (238,486174 miljonit võtit/s)

0,69
18,690048 ms (224,413765 miljonit võtit/s)

0,72
20,278816 ms (206,831789 miljonit võtit/s)

0,75
22,545408 ms (186,038058 miljonit võtit/s)

0,78
26,053312 ms (160,989275 miljonit võtit/s)

0,81
31,895008 ms (131,503463 miljonit võtit/s)

0,84
42,103294 ms (99,619378 miljonit võtit/s)

0,87
61,849056 ms (67,815164 miljonit võtit/s)

0,90
105,695999 ms (39,682713 miljonit võtit/s)

0,94
240,204636 ms (17,461378 miljonit võtit/s)

Kuna kasutuse koefitsient tõuseb, väheneb tulemuslikkus. See pole enamikul juhtudel soovitav. Kui rakendus sisestab tabelisse elemente ja seejärel need kõrvaldada (näiteks raamatu sõnade loendamisel), siis see ei ole probleem. Kuid kui rakendus kasutab pikaajalist rippmenahtla (nt graafikutoimetajas mitte tühi osa piltide hoidmiseks, kui kasutaja sageli sisestab ja eemaldab teavet), võib selline käitumine tekitada ebamugavusi.

Ja ma mõõtsin rippmenahtla sondimis sügavust pärast 64 miljoni sisestamist (kasutuse koefitsient 0,5). Keskmine sügavus oli 0,4774, nii et enamik võtmeid asus kas parimas võimalikus slotis või ühe sloti kaugusel parimast asukohast. Maksimaalne sondimissügavus oli 60.

Siis mõõtsin sondimissügavust tabelis, kus oli 124 miljonit sisestamist (kasutuse koefitsient 0,97). Keskmine sügavus oli juba 10,1757, maksimaalne — 6474 (!!). Lineaarse sondimise tulemuslikkus langeb märkimisväärselt, kui kasutuse koefitsient on kõrge.

Parima on selle hash-tabeli madal kasutussuhte säilitamine. Kuid see tähendaks, et me suurendame jõudlust mälutarbimise arvelt. Õnneks on see 32-bitiste võtmete ja väärtuste korral põhjendatud. Kui ülaltoodud näites, kus tabelis on 128 miljonit elementi, hoida kasutussuhteks 0,25, saame me tabelisse mahutada mitte rohkem kui 32 miljonit elementi, samas kui ülejäänud 96 miljonit slotti lähevad kaotsi — iga paari jaoks 8 baiti, kokku 768 MB kaotatud mälu.

Pange tähele, et jutt on videokaardi mälu kaotusest, mis on väärtuslikum ressurss kui süsteemimälu. Kuigi enamikul kaasaegsetest lauaarvutite videokaardidest, mis toetavad CUDA-d, on vähemalt 4 GB mälu (artikli kirjutamise ajal on NVIDIA 2080 Ti-l 11 GB), suurte mahtude kaotamine ei ole siiski kõige mõistlikum lahendus.

Hiljem kirjutan ma lähemalt videokaartide jaoks hash-tabelite loomisest, millel ei esine sügavusprobleeme, samuti uuesti kasutatavate eemaldatud slotide meetoditest.

Sügavuse mõõtmine

Selleks, et määrata võtme sügavus, saame välja võtta võtme hash'i (selle ideaalne indeks tabelis) tegelikust tabeli indeksist:

// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);

Kuna kaht kahendnumbrite maagia ja see, et hash-tabeli maht on kahe astmest, toimib see lähenemine isegi siis, kui võtme indeks kantakse tabeli algusesse. Võtame võtme, mis hashitakse 1, kuid on sisestatud slotti 3. Siis tabeli maht 4 puhul saame (3 — 1) & 3, mis on võrreldav 2.

Kokkuvõte

Kui teil on küsimusi või kommentaare, kirjutage mulle Twitteris või avage uus teema hoidlad.

See kood on inspireeritud suurepärasest artiklist:

• Maailma lihtsaim lukustamata hash-tabel
• Bloqueeringu vaba ootevaba hash-tabel

Tulevikus jätkan ma videokaartidele mõeldud hash-tabelite rakenduste kirjutamist ja analüüsin nende jõudlust. Plaanide hulgas on ahelal põhinev hashimine, Robin Hoodi hashimine ning kukeharu hashimine, kasutades aatomoperatsioone andmestruktuurides, mis sobivad videokaartidele.

Allikas: habr.com