Julkaisin sen Githubissa .
Se on yksinkertainen GPU-hajautustaulukko, joka pystyy käsittelemään satoja miljoonia lisäyksiä sekunnissa. NVIDIA GTX 1060 -kannettavassani koodi lisää 64 miljoonaa satunnaisesti luotua avainarvoparia noin 210 ms:ssa ja poistaa 32 miljoonaa paria noin 64 ms:ssa.
Toisin sanoen kannettavan tietokoneen nopeus on noin 300 miljoonaa lisäystä/s ja 500 miljoonaa poistoa/s.
Taulukko on kirjoitettu CUDA-kielellä, vaikka samaa tekniikkaa voidaan soveltaa HLSL:ään tai GLSL:ään. Toteutuksessa on useita rajoituksia näytönohjaimen korkean suorituskyvyn varmistamiseksi:
- Vain 32-bittiset avaimet ja samat arvot käsitellään.
- Hash-taulukossa on kiinteä koko.
- Ja tämän koon on oltava yhtä suuri kuin kaksi tehoa.
Avaimia ja arvoja varten sinun on varattava yksinkertainen erotinmerkki (yllä olevassa koodissa tämä on 0xffffffff).
Hash-pöytä ilman lukkoja
Hajautustaulukko käyttää avointa osoitusta kanssa , eli se on yksinkertaisesti joukko avainarvopareja, jotka on tallennettu muistiin ja joilla on ylivoimainen välimuistin suorituskyky. Samaa ei voida sanoa ketjuttamisesta, joka sisältää osoittimen etsimisen linkitetystä luettelosta. Hash-taulukko on yksinkertainen taulukko, joka tallentaa elementtejä KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Taulukon koko on kahden potenssi, ei alkuluku, koska yksi nopea käsky riittää pow2/AND-maskin käyttöön, mutta moduulioperaattori on paljon hitaampi. Tämä on tärkeää lineaarisen mittauksen tapauksessa, koska lineaaritaulukon haussa aika-indeksi on käärittävä jokaiseen paikkaan. Tämän seurauksena toiminnan kustannukset lisätään modulo jokaiseen paikkaan.
Taulukko tallentaa vain kunkin elementin avaimen ja arvon, ei avaimen tiivistettä. Koska taulukko tallentaa vain 32-bittiset avaimet, hash lasketaan erittäin nopeasti. Yllä oleva koodi käyttää Murmur3-tiivistettä, joka suorittaa vain muutaman siirron, XOR:n ja kertolaskujen.
Hajautustaulukko käyttää lukitussuojaustekniikoita, jotka ovat riippumattomia muistijärjestyksestä. Vaikka jotkin kirjoitustoiminnot häiritsevät muiden tällaisten toimintojen järjestystä, hash-taulukko säilyttää silti oikean tilan. Puhumme tästä alla. Tekniikka toimii hyvin näytönohjainkorteilla, jotka ajavat tuhansia säikeitä samanaikaisesti.
Hajautustaulukon avaimet ja arvot alustetaan tyhjiksi.
Koodia voidaan muokata käsittelemään myös 64-bittisiä avaimia ja arvoja. Avaimet vaativat atomiluku-, kirjoitus- ja vertailu- ja vaihtotoimintoja. Ja arvot vaativat atomiluku- ja kirjoitusoperaatioita. Onneksi CUDA:ssa 32- ja 64-bittisten arvojen luku-kirjoitustoiminnot ovat atomisia niin kauan kuin ne ovat luonnollisesti kohdistettuja (katso alla). ), ja nykyaikaiset näytönohjaimet tukevat 64-bittisiä atomisia vertailu- ja vaihtotoimintoja. Tietenkin, kun siirrytään 64-bittiseen, suorituskyky heikkenee hieman.
Hash-taulukon tila
Jokaisella hash-taulukon avainarvoparilla voi olla yksi neljästä tilasta:
- Avain ja arvo ovat tyhjiä. Tässä tilassa hash-taulukko alustetaan.
- Avain on kirjoitettu, mutta arvoa ei ole vielä kirjoitettu. Jos toinen säie lukee parhaillaan tietoja, se palaa tyhjänä. Tämä on normaalia, sama asia olisi tapahtunut, jos toinen suoritussäie olisi toiminut hieman aikaisemmin, ja puhumme samanaikaisesta tietorakenteesta.
- Sekä avain että arvo tallennetaan.
- Arvo on muiden suoritussäikeiden käytettävissä, mutta avain ei ole vielä. Tämä voi tapahtua, koska CUDA-ohjelmointimallissa on löyhästi järjestetty muistimalli. Tämä on normaalia; joka tapauksessa avain on edelleen tyhjä, vaikka arvo ei enää olisikaan.
Tärkeä vivahde on, että kun avain on kirjoitettu paikkaan, se ei enää liiku - vaikka avain poistetaan, puhumme tästä alla.
Hajautustaulukkokoodi toimii jopa löyhästi järjestetyissä muistimalleissa, joissa muistin luku- ja kirjoitusjärjestystä ei tunneta. Kun tarkastelemme lisäystä, hakua ja poistamista hash-taulukossa, muista, että jokainen avain-arvo-pari on jossakin neljästä yllä kuvatusta tilasta.
Lisääminen hash-taulukkoon
CUDA-funktio, joka lisää avainarvopareja hash-taulukkoon, näyttää tältä:
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);
}
}
Avaimen lisäämiseksi koodi iteroidaan hash-taulukkotaulukon läpi alkaen lisätyn avaimen hashista. Jokainen taulukon paikka suorittaa atomivertailu- ja vaihtooperaation, joka vertaa kyseisen paikan avainta tyhjään. Jos yhteensopimattomuus havaitaan, paikan avain päivitetään lisätyllä avaimella ja sitten palautetaan alkuperäinen paikan avain. Jos tämä alkuperäinen avain oli tyhjä tai vastasi lisättyä avainta, koodi löysi sopivan paikan lisättäväksi ja lisäsi lisätyn arvon aukkoon.
Jos yhdessä ytimen kutsussa gpu_hashtable_insert() samalla avaimella on useita elementtejä, minkä tahansa niiden arvot voidaan kirjoittaa avainpaikkaan. Tätä pidetään normaalina: yksi avainarvon kirjoituksista puhelun aikana onnistuu, mutta koska tämä kaikki tapahtuu rinnakkain useissa suoritussäikeissä, emme voi ennustaa, mikä muistikirjoitus on viimeinen.
Hash-taulukon haku
Koodi avainten hakua varten:
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);
}
}
Löytääksemme taulukkoon tallennetun avaimen arvon iteroimme taulukon läpi alkaen etsimämme avaimen hajautusarvosta. Jokaisessa paikassa tarkistamme, onko avain se, jota etsimme, ja jos on, palautamme sen arvon. Tarkistamme myös, onko avain tyhjä, ja jos on, keskeytämme haun.
Jos emme löydä avainta, koodi palauttaa tyhjän arvon.
Kaikki nämä hakutoiminnot voidaan suorittaa samanaikaisesti lisäämällä ja poistamalla. Jokaisella taulukon parilla on yksi neljästä edellä kuvatusta virran tilasta.
Poistaminen hash-taulukossa
Koodi avainten poistamiseen:
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);
}
}
Avaimen poistaminen tapahtuu epätavallisella tavalla: jätämme avaimen taulukkoon ja merkitsemme sen arvon (ei itse avainta) tyhjäksi. Tämä koodi on hyvin samanlainen lookup(), paitsi että kun avaimesta löytyy osuma, se tekee sen arvon tyhjäksi.
Kuten edellä mainittiin, kun avain on kirjoitettu paikkaan, sitä ei enää siirretä. Vaikka elementti poistetaan taulukosta, avain pysyy paikallaan, sen arvo yksinkertaisesti tyhjenee. Tämä tarkoittaa, että meidän ei tarvitse käyttää atomin kirjoitusoperaatiota paikan arvolle, koska sillä ei ole väliä, onko nykyinen arvo tyhjä vai ei - se tulee silti tyhjäksi.
Hajautustaulukon koon muuttaminen
Voit muuttaa hash-taulukon kokoa luomalla suuremman taulukon ja lisäämällä siihen ei-tyhjiä elementtejä vanhasta taulukosta. En ottanut tätä toimintoa käyttöön, koska halusin pitää mallikoodin yksinkertaisena. Lisäksi CUDA-ohjelmissa muistin varaus tehdään usein isäntäkoodissa CUDA-ytimen sijaan.
Artikkelissa kuvataan, kuinka tällaista lukitussuojattua tietorakennetta muutetaan.
Kilpailukyky
Yllä olevissa toimintokoodinpätkissä gpu_hashtable_insert(), _lookup() и _delete() käsitellä yhtä avainarvoparia kerrallaan. Ja alemmas gpu_hashtable_insert(), _lookup() и _delete() käsittelemään useita pareja rinnakkain, jokainen pari erillisessä GPU-suoritussäikeessä:
// 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);
}
}
Lukitusta kestävä hash-taulukko tukee samanaikaisia lisäyksiä, hakuja ja poistoja. Koska avain-arvo-parit ovat aina jossakin neljästä tilasta ja avaimet eivät liiku, taulukko takaa oikeellisuuden myös silloin, kun eri tyyppisiä operaatioita käytetään samanaikaisesti.
Jos kuitenkin käsittelemme lisäys- ja poistoerän rinnakkain ja jos syöteparien ryhmä sisältää päällekkäisiä avaimia, emme voi ennustaa, mitkä parit "voittavat" - kirjoitetaan hash-taulukkoon viimeisenä. Oletetaan, että kutsuimme lisäyskoodia syöteparien joukolla A/0 B/1 A/2 C/3 A/4. Kun koodi on valmis, paritetaan B/1 и C/3 ovat taatusti läsnä taulukossa, mutta samaan aikaan mikä tahansa pareista näkyy siinä A/0, A/2 tai A/4. Tämä voi olla ongelma tai ei - kaikki riippuu sovelluksesta. Saatat tietää etukäteen, että syöttötaulukossa ei ole päällekkäisiä avaimia, tai et ehkä välitä, mikä arvo kirjoitettiin viimeksi.
Jos tämä on sinulle ongelma, sinun on erotettava kaksoisparit eri CUDA-järjestelmäkutsuiksi. CUDA:ssa kaikki ydintä kutsuvat toiminnot valmistuvat aina ennen seuraavaa ytimen kutsua (ainakin yhden säikeen sisällä. Eri säikeissä ytimet suoritetaan rinnakkain). Yllä olevassa esimerkissä, jos kutsut yhtä ydintä käyttämällä A/0 B/1 A/2 C/3, ja toinen kanssa A/4, sitten avain A saa arvon 4.
Nyt puhutaan siitä, pitäisikö funktioiden olla lookup() и delete() käytä tavallista tai haihtuvaa osoitinta hajautustaulukon parien joukkoon. toteaa, että:
Kääntäjä voi optimoida lukemisen ja kirjoittamisen globaaliin tai jaettuun muistiin... Nämä optimoinnit voidaan poistaa käytöstä avainsanalla
volatile: ... kaikki viittaukset tähän muuttujaan käännetään todelliseksi muistin luku- tai kirjoituskäskyksi.
Oikeudenmukaisuusnäkökohdat eivät vaadi soveltamista volatile. Jos suoritussäie käyttää välimuistissa olevaa arvoa aikaisemmasta lukutoiminnosta, se käyttää hieman vanhentuneita tietoja. Mutta silti, tämä on tietoa hajautustaulukon oikeasta tilasta tietyllä ytimen kutsun hetkellä. Jos haluat käyttää uusimpia tietoja, voit käyttää hakemistoa volatile, mutta silloin suorituskyky heikkenee hieman: testieni mukaan 32 miljoonaa elementtiä poistettaessa nopeus laski 500 miljoonasta poistosta sekunnissa 450 miljoonaan poistoon sekunnissa.
Suorituskyky
Testissä 64 miljoonan elementin lisäämisestä ja niistä 32 miljoonan poistamisesta kilpailee välillä std::unordered_map ja grafiikkasuorittimelle ei ole käytännössä yhtään hash-taulukkoa:
std::unordered_map käytti 70 691 ms elementtien lisäämiseen ja poistamiseen ja niiden vapauttamiseen unordered_map (miljoonista elementeistä eroon pääseminen vie paljon aikaa, koska sisällä unordered_map useita muistivarauksia tehdään). Rehellisesti sanottuna, std:unordered_map täysin erilaisia rajoituksia. Se on yksi suoritinsäie, joka tukee kaikenkokoisia avainarvoja, toimii hyvin korkeilla käyttöasteikoilla ja näyttää vakaan suorituskyvyn useiden poistojen jälkeen.
GPU:n ja ohjelmien välisen viestinnän hash-taulukon kesto oli 984 ms. Tähän sisältyy aika, joka kuluu taulukon sijoittamiseen muistiin ja poistamiseen (1 Gt muistin varaaminen kerralla, mikä vie jonkin aikaa CUDA:ssa), elementtien lisäämiseen ja poistamiseen sekä niiden iterointiin. Myös kaikki kopiot näytönohjaimen muistista ja muistista otetaan huomioon.
Itse hash-taulukon valmistuminen kesti 271 ms. Tämä sisältää ajan, jonka näytönohjain käyttää elementtien asettamiseen ja poistamiseen, eikä siinä oteta huomioon aikaa, joka kuluu muistiin kopiointiin ja tuloksena olevan taulukon iterointiin. Jos GPU-taulukko elää pitkään tai jos hash-taulukko on kokonaan näytönohjaimen muistissa (esimerkiksi tiivistetaulukon luomiseksi, jota käytetään muussa GPU-koodissa, ei keskusprosessorissa), testitulos on relevantti.
Näytönohjaimen hash-taulukko osoittaa korkean suorituskyvyn suuren suorituskyvyn ja aktiivisen rinnakkaistoiminnon ansiosta.
Rajoitukset
Hash-taulukon arkkitehtuurissa on muutama huomioitava ongelma:
- Lineaarista mittaamista vaikeuttaa klusterointi, jolloin taulukon avaimet eivät sijoittu täydellisesti.
- Näppäimiä ei poisteta toimintoa käyttämällä
deleteja ajan myötä ne sotkevat pöydän.
Tämän seurauksena hash-taulukon suorituskyky voi vähitellen heiketä, varsinkin jos se on olemassa pitkään ja siinä on useita lisäyksiä ja poistoja. Yksi tapa lieventää näitä haittoja on tiivistää uudelleen uuteen taulukkoon, jonka käyttöaste on melko alhainen, ja suodattaa poistetut avaimet pois uudelleenhajauksen aikana.
Kuvattujen ongelmien havainnollistamiseksi luon yllä olevan koodin avulla taulukon, jossa on 128 miljoonaa elementtiä, ja kierrän 4 miljoonan elementin läpi, kunnes olen täyttänyt 124 miljoonaa paikkaa (käyttöaste noin 0,96). Tässä on tulostaulukko, jokainen rivi on CUDA-ytimen kutsu 4 miljoonan uuden elementin lisäämiseksi yhteen hash-taulukkoon:
Käyttöaste
Lisäyksen kesto 4 194 304 elementtiä
0,00
11,608448 ms (361,314798 miljoonaa näppäintä/sek.)
0,03
11,751424 ms (356,918799 miljoonaa näppäintä/sek.)
0,06
11,942592 ms (351,205515 miljoonaa näppäintä/sek.)
0,09
12,081120 ms (347,178429 miljoonaa näppäintä/sek.)
0,12
12,242560 ms (342,600233 miljoonaa näppäintä/sek.)
0,16
12,396448 ms (338,347235 miljoonaa näppäintä/sek.)
0,19
12,533024 ms (334,660176 miljoonaa näppäintä/sek.)
0,22
12,703328 ms (330,173626 miljoonaa näppäintä/sek.)
0,25
12,884512 ms (325,530693 miljoonaa näppäintä/sek.)
0,28
13,033472 ms (321,810182 miljoonaa näppäintä/sek.)
0,31
13,239296 ms (316,807174 miljoonaa näppäintä/sek.)
0,34
13,392448 ms (313,184256 miljoonaa näppäintä/sek.)
0,37
13,624000 ms (307,861434 miljoonaa näppäintä/sek.)
0,41
13,875520 ms (302,280855 miljoonaa näppäintä/sek.)
0,44
14,126528 ms (296,909756 miljoonaa näppäintä/sek.)
0,47
14,399328 ms (291,284699 miljoonaa näppäintä/sek.)
0,50
14,690304 ms (285,515123 miljoonaa näppäintä/sek.)
0,53
15,039136 ms (278,892623 miljoonaa näppäintä/sek.)
0,56
15,478656 ms (270,973402 miljoonaa näppäintä/sek.)
0,59
15,985664 ms (262,379092 miljoonaa näppäintä/sek.)
0,62
16,668673 ms (251,627968 miljoonaa näppäintä/sek.)
0,66
17,587200 ms (238,486174 miljoonaa näppäintä/sek.)
0,69
18,690048 ms (224,413765 miljoonaa näppäintä/sek.)
0,72
20,278816 ms (206,831789 miljoonaa näppäintä/sek.)
0,75
22,545408 ms (186,038058 miljoonaa näppäintä/sek.)
0,78
26,053312 ms (160,989275 miljoonaa näppäintä/sek.)
0,81
31,895008 ms (131,503463 miljoonaa näppäintä/sek.)
0,84
42,103294 ms (99,619378 miljoonaa näppäintä/sek.)
0,87
61,849056 ms (67,815164 miljoonaa näppäintä/sek.)
0,90
105,695999 ms (39,682713 miljoonaa näppäintä/sek.)
0,94
240,204636 ms (17,461378 miljoonaa näppäintä/sek.)
Kun käyttöaste kasvaa, suorituskyky heikkenee. Tämä ei ole toivottavaa useimmissa tapauksissa. Jos sovellus lisää elementtejä taulukkoon ja sitten hylkää ne (esimerkiksi laskettaessa sanoja kirjassa), tämä ei ole ongelma. Mutta jos sovellus käyttää pitkäikäistä hash-taulukkoa (esimerkiksi grafiikkaeditorissa kuvien ei-tyhjien osien tallentamiseen, joihin käyttäjä usein lisää ja poistaa tietoja), tämä toiminta voi olla ongelmallista.
Ja mittasi hash-taulukon mittaussyvyyden 64 miljoonan lisäyksen jälkeen (käyttökerroin 0,5). Keskisyvyys oli 0,4774, joten useimmat näppäimet olivat joko parhaassa mahdollisessa paikassa tai yhden paikan päässä parhaasta asennosta. Suurin luotaussyvyys oli 60.
Sitten mittasin mittaussyvyyden pöydällä, jossa oli 124 miljoonaa terää (käyttökerroin 0,97). Keskisyvyys oli jo 10,1757 ja suurin - 6474 (!!). Lineaarisen tunnistuksen suorituskyky heikkenee merkittävästi korkeilla käyttöasteilla.
On parasta pitää tämän hash-taulukon käyttöaste alhaisena. Mutta sitten lisäämme suorituskykyä muistin kulutuksen kustannuksella. Onneksi 32-bittisten avainten ja arvojen tapauksessa tämä voi olla perusteltua. Jos yllä olevassa esimerkissä taulukossa, jossa on 128 miljoonaa elementtiä, säilytetään käyttökerroin 0,25, voimme sijoittaa siihen enintään 32 miljoonaa elementtiä ja loput 96 miljoonaa paikkaa menetetään - 8 tavua jokaista paria kohti , 768 Mt menetettyä muistia.
Huomaa, että puhumme näytönohjaimen muistin menetyksestä, joka on arvokkaampi resurssi kuin järjestelmämuisti. Vaikka useimmissa nykyaikaisissa CUDAa tukevissa työpöytänäytönohjainkorteissa on vähintään 4 Gt muistia (kirjoitushetkellä NVIDIA 2080 Ti:ssä on 11 Gt), ei silti olisi viisain päätös menettää tällaisia määriä.
Myöhemmin kirjoitan lisää hash-taulukoiden luomisesta näytönohjaimille, joilla ei ole ongelmia mittaussyvyyden kanssa, sekä tavoista käyttää poistettuja paikkoja uudelleen.
Luotainsyvyyden mittaus
Avaimen mittaussyvyyden määrittämiseksi voimme poimia avaimen tiivisteen (sen ihanteellinen taulukkoindeksi) sen varsinaisesta taulukkoindeksistä:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Kahden kahden komplementin binäärilukujen taikuudesta ja siitä, että hash-taulukon kapasiteetti on kaksi potenssilla kaksi, tämä lähestymistapa toimii myös silloin, kun avainindeksi siirretään taulukon alkuun. Otetaan avain, joka on hajautettu arvoon 1, mutta joka on lisätty paikkaan 3. Sitten taulukolle, jonka kapasiteetti on 4, saamme (3 — 1) & 3, joka vastaa 2.
Johtopäätös
Jos sinulla on kysyttävää tai kommentteja, lähetä minulle sähköpostia osoitteeseen tai avaa uusi aihe.
Tämä koodi on kirjoitettu erinomaisten artikkelien inspiraation alaisena:
Tulevaisuudessa jatkan kirjoittamista näytönohjainten hash-taulukoiden toteutuksista ja analysoin niiden suorituskykyä. Suunnitelmiini kuuluu ketjuttaminen, Robin Hood -hajautus ja käkihajautus käyttämällä atomioperaatioita tietorakenteissa, jotka ovat GPU-ystävällisiä.
Lähde: will.com