Paskelbiau jį Github .
Tai paprasta GPU maišos lentelė, galinti apdoroti šimtus milijonų įterpimų per sekundę. Mano NVIDIA GTX 1060 nešiojamame kompiuteryje kodas įterpia 64 milijonus atsitiktinai sugeneruotų raktų verčių porų maždaug per 210 ms ir pašalina 32 milijonus porų per maždaug 64 ms.
Tai reiškia, kad nešiojamojo kompiuterio greitis yra maždaug 300 milijonų įterpimų per sekundę ir 500 milijonų ištrynimų per sekundę.
Lentelė parašyta CUDA, nors tą pačią techniką galima pritaikyti HLSL arba GLSL. Diegimas turi keletą apribojimų, užtikrinančių aukštą vaizdo plokštės našumą:
- Apdorojami tik 32 bitų raktai ir tos pačios reikšmės.
- Maišos lentelės dydis yra fiksuotas.
- Ir šis dydis turi būti lygus dviem galiai.
Raktams ir reikšmėms turite rezervuoti paprastą skyriklio žymeklį (aukščiau pateiktame kode tai yra 0xffffffff).
Maišos stalas be spynų
Maišos lentelėje naudojamas atviras adresas su , tai yra, tai tiesiog raktų ir reikšmių porų masyvas, saugomas atmintyje ir pasižymintis puikiu talpyklos našumu. To negalima pasakyti apie grandinės sujungimą, kai reikia ieškoti rodyklės susietame sąraše. Maišos lentelė yra paprastas masyvas, kuriame saugomi elementai KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Lentelės dydis yra dviejų laipsnis, o ne pirminis skaičius, nes pow2/AND kaukei pritaikyti užtenka vienos greitos instrukcijos, tačiau modulio operatorius yra daug lėtesnis. Tai svarbu linijinio zondavimo atveju, nes tiesinės lentelės peržvalgoje lizdo indeksas turi būti įtrauktas į kiekvieną lizdą. Dėl to operacijos kaina pridedama modulio kiekviename lizde.
Lentelėje saugomas tik kiekvieno elemento raktas ir reikšmė, o ne rakto maiša. Kadangi lentelėje saugomi tik 32 bitų raktai, maiša apskaičiuojama labai greitai. Aukščiau pateiktame kode naudojama Murmur3 maiša, kuri atlieka tik keletą poslinkių, XOR ir daugybos.
Maišos lentelėje naudojami užrakinimo apsaugos būdai, kurie nepriklauso nuo atminties tvarkos. Net jei kai kurios rašymo operacijos sutrikdo kitų tokių operacijų tvarką, maišos lentelė vis tiek išlaikys teisingą būseną. Apie tai kalbėsime žemiau. Ši technika puikiai veikia su vaizdo plokštėmis, kuriose vienu metu veikia tūkstančiai gijų.
Raktai ir reikšmės maišos lentelėje inicijuojamos į tuščius.
Kodą galima modifikuoti, kad būtų galima apdoroti ir 64 bitų raktus bei reikšmes. Klavišams reikalingos atominės skaitymo, rašymo ir palyginimo bei keitimo operacijos. O vertybėms reikalingos atominės skaitymo ir rašymo operacijos. Laimei, CUDA 32 ir 64 bitų reikšmių skaitymo ir rašymo operacijos yra atominės tol, kol jos yra natūraliai suderintos (žr. toliau). ), o šiuolaikinės vaizdo plokštės palaiko 64 bitų atomines palyginimo ir keitimo operacijas. Žinoma, pereinant prie 64 bitų našumas šiek tiek sumažės.
Maišos lentelės būsena
Kiekviena maišos lentelės rakto ir verčių pora gali turėti vieną iš keturių būsenų:
- Raktas ir reikšmė tušti. Šioje būsenoje maišos lentelė inicijuojama.
- Raktas užrašytas, bet vertė dar neparašyta. Jei kita gija šiuo metu skaito duomenis, ji grįžta tuščia. Tai normalu, tas pats būtų nutikę, jei kita vykdymo gija būtų veikusi šiek tiek anksčiau, o mes kalbame apie lygiagrečią duomenų struktūrą.
- Įrašomi ir raktas, ir vertė.
- Vertė prieinama kitoms vykdymo gijomis, bet rakto dar nėra. Taip gali atsitikti, nes CUDA programavimo modelis turi laisvai sutvarkytą atminties modelį. Tai normalu; bet kuriuo atveju raktas vis dar tuščias, net jei reikšmės nebėra.
Svarbus niuansas yra tai, kad kai raktas buvo įrašytas į lizdą, jis nebejuda – net jei raktas ištrintas, apie tai kalbėsime žemiau.
Maišos lentelės kodas netgi veikia su laisvai išdėstytais atminties modeliais, kuriuose atminties skaitymo ir rašymo tvarka nežinoma. Žvelgdami į įterpimą, paiešką ir ištrynimą maišos lentelėje, atminkite, kad kiekviena rakto ir verčių pora yra vienoje iš keturių aukščiau aprašytų būsenų.
Įterpimas į maišos lentelę
CUDA funkcija, kuri įterpia raktų ir reikšmių poras į maišos lentelę, atrodo taip:
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);
}
}
Norint įterpti raktą, kodas kartojamas maišos lentelės masyvu, pradedant nuo įterpto rakto maišos. Kiekvienas masyvo lizdas atlieka atominę palyginimo ir keitimo operaciją, kuri lygina raktą toje vietoje su tuščiu. Jei aptinkamas neatitikimas, lizde esantis raktas atnaujinamas įdėtu raktu, o tada grąžinamas pradinis lizdo raktas. Jei šis originalus raktas buvo tuščias arba atitiko įdėtą raktą, tada kodas rado tinkamą įterpimui skirtą lizdą ir įterptą reikšmę įterpė į lizdą.
Jei per vieną branduolio skambutį gpu_hashtable_insert() yra keli elementai su tuo pačiu raktu, tada bet kurią jų reikšmę galima įrašyti į rakto lizdą. Tai laikoma normalia: vienas iš rakto-reikšmių įrašų skambučio metu bus sėkmingas, bet kadangi visa tai vyksta lygiagrečiai keliose vykdymo gijose, negalime numatyti, kuris atminties įrašymas bus paskutinis.
Maišos lentelės paieška
Raktų paieškos kodas:
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);
}
}
Norėdami rasti lentelėje saugomo rakto reikšmę, kartojame masyvą, pradedant nuo ieškomo rakto maišos. Kiekviename lizde patikriname, ar raktas yra tas, kurio ieškome, o jei taip, grąžiname jo vertę. Taip pat patikriname, ar raktas tuščias, o jei taip, paiešką nutraukiame.
Jei negalime rasti rakto, kodas grąžina tuščią reikšmę.
Visos šios paieškos operacijos gali būti atliekamos vienu metu įterpiant ir ištrinant. Kiekviena lentelės pora turės vieną iš keturių aukščiau aprašytų srauto būsenų.
Ištrynimas maišos lentelėje
Raktų ištrynimo kodas:
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);
}
}
Rakto ištrynimas atliekamas neįprastu būdu: paliekame raktą lentelėje ir pažymime jo reikšmę (ne patį raktą) kaip tuščią. Šis kodas labai panašus į lookup(), išskyrus tai, kad kai ant rakto randama atitiktis, jo reikšmė tuščia.
Kaip minėta aukščiau, kai raktas įrašomas į lizdą, jis nebejudinamas. Net ištrinus elementą iš lentelės, raktas lieka vietoje, jo reikšmė tiesiog tampa tuščia. Tai reiškia, kad lizdo reikšmei nereikia naudoti atominės rašymo operacijos, nes nesvarbu, ar dabartinė reikšmė tuščia, ar ne – ji vis tiek taps tuščia.
Maišos lentelės dydžio keitimas
Maišos lentelės dydį galite pakeisti sukurdami didesnę lentelę ir į ją įterpdami netuščius elementus iš senos lentelės. Šios funkcijos neįdiegiau, nes norėjau, kad pavyzdinis kodas būtų paprastas. Be to, CUDA programose atminties paskirstymas dažnai atliekamas pagrindinio kompiuterio kode, o ne CUDA branduolyje.
Straipsnis aprašoma, kaip modifikuoti tokią nuo užrakto apsaugotą duomenų struktūrą.
Konkurencingumas
Aukščiau pateiktuose funkcijos kodo fragmentuose gpu_hashtable_insert(), _lookup() и _delete() vienu metu apdoroti vieną rakto-reikšmių porą. Ir žemesnė gpu_hashtable_insert(), _lookup() и _delete() lygiagrečiai apdoroti porų masyvą, kiekviena pora atskiroje GPU vykdymo gijoje:
// 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);
}
}
Neužrakinama maišos lentelė palaiko vienu metu atliekamus įterpimus, paieškas ir trynimus. Kadangi raktų ir reikšmių poros visada yra vienoje iš keturių būsenų ir raktai nejuda, lentelė garantuoja teisingumą net tada, kai vienu metu naudojamos skirtingų tipų operacijos.
Tačiau jei lygiagrečiai apdorosime įterpimų ir ištrynimų paketą, o įvesties porų masyve yra pasikartojantys raktai, negalėsime nuspėti, kurios poros „laimės“ – bus įrašytos į maišos lentelę paskutinės. Tarkime, kad įterpimo kodą pavadinome porų įvesties masyvu A/0 B/1 A/2 C/3 A/4. Kai kodas baigtas, suporuojama B/1 и C/3 garantuotai bus lentelėje, bet tuo pačiu metu joje atsiras bet kuri iš porų A/0, A/2 arba A/4. Tai gali būti problema arba ne – viskas priklauso nuo programos. Galite iš anksto žinoti, kad įvesties masyve nėra pasikartojančių raktų arba jums gali nerūpėti, kuri reikšmė buvo įrašyta paskutinė.
Jei jums tai kelia problemų, tuomet turite atskirti pasikartojančias poras į skirtingus CUDA sistemos skambučius. CUDA bet kuri branduolį iškviečianti operacija visada baigiama prieš kitą branduolio iškvietimą (bent jau vienoje gijoje. Skirtingose gijose branduoliai vykdomi lygiagrečiai). Aukščiau pateiktame pavyzdyje, jei iškviečiate vieną branduolį su A/0 B/1 A/2 C/3, o kitas su A/4, tada raktas A gaus vertę 4.
Dabar pakalbėkime apie tai, ar funkcijos turėtų būti lookup() и delete() naudokite paprastą arba nepastovią rodyklę į porų masyvą maišos lentelėje. Teigia, kad:
Kompiliatorius gali pasirinkti optimizuoti skaitymą ir rašymą į visuotinę arba bendrąją atmintį... Šiuos optimizavimus galima išjungti naudojant raktinį žodį
volatile: ... bet kokia nuoroda į šį kintamąjį sukompiliuojama į realią atminties skaitymo arba rašymo instrukciją.
Teisingumo sumetimais nereikia taikyti volatile. Jei vykdymo gija naudoja talpykloje saugomą reikšmę iš ankstesnės skaitymo operacijos, tada ji naudos šiek tiek pasenusią informaciją. Bet vis tiek tai yra informacija iš teisingos maišos lentelės būsenos tam tikru branduolio iškvietimo momentu. Jei reikia naudoti naujausią informaciją, galite naudoti rodyklę volatile, bet tada našumas šiek tiek sumažės: pagal mano testus, ištrinant 32 milijonus elementų, greitis sumažėjo nuo 500 milijonų trynimų/sek iki 450 milijonų trynimų/sek.
Našumas
Atliekant bandymą įterpti 64 milijonus elementų ir ištrinti 32 milijonus iš jų, konkurencija tarp std::unordered_map ir praktiškai nėra GPU maišos lentelės:
std::unordered_map praleido 70 691 ms įdėdamas ir išimdamas elementus, o paskui juos atlaisvindamas unordered_map (atsikratymas iš milijonų elementų užima daug laiko, nes viduje unordered_map atliekami keli atminties paskirstymai). Atvirai kalbant, std:unordered_map visiškai skirtingi apribojimai. Tai viena CPU vykdymo gija, palaiko bet kokio dydžio raktų reikšmes, gerai veikia esant dideliam panaudojimo lygiui ir rodo stabilų veikimą po kelių trynimų.
GPU ir tarpprograminio ryšio maišos lentelės trukmė buvo 984 ms. Tai apima laiką, praleistą lentelės įdėjimui į atmintį ir jos ištrynimui (vieną kartą skiriant 1 GB atminties, o tai užtrunka šiek tiek laiko CUDA), elementų įterpimui ir trynimui bei kartojimui per juos. Taip pat atsižvelgiama į visas kopijas į vaizdo plokštės atmintį ir iš jos.
Pati maišos lentelė užtruko 271 ms. Tai apima laiką, kurį vaizdo plokštė praleidžia įdedant ir ištrinant elementus, ir neatsižvelgiama į laiką, praleistą kopijuojant į atmintį ir kartojant gautą lentelę. Jei GPU lentelė veikia ilgą laiką arba jei maišos lentelė yra visiškai vaizdo plokštės atmintyje (pavyzdžiui, norint sukurti maišos lentelę, kurią naudos kitas GPU kodas, o ne centrinis procesorius), tada testo rezultatas yra aktualus.
Vaizdo plokštės maišos lentelė demonstruoja didelį našumą dėl didelio pralaidumo ir aktyvaus lygiagretinimo.
Trūkumai
Maišos lentelės architektūra turi keletą problemų, kurias reikia žinoti:
- Linijinį zondavimą apsunkina grupavimas, dėl kurio klavišai lentelėje yra išdėstyti prasčiau nei idealiai.
- Naudojant šią funkciją klavišai nepašalinami
deleteir laikui bėgant jie užgriozdina stalą.
Dėl to maišos lentelės našumas gali palaipsniui pablogėti, ypač jei ji egzistuoja ilgą laiką ir turi daug įterpimų ir ištrynimų. Vienas iš būdų sušvelninti šiuos trūkumus yra perrašyti į naują lentelę su gana žemu panaudojimo rodikliu ir išfiltruoti pašalintus raktus atliekant pakartotinę maišą.
Siekdamas iliustruoti aprašytas problemas, naudosiu aukščiau pateiktą kodą, kad sukurčiau lentelę su 128 milijonais elementų ir peržiūrėsiu 4 milijonus elementų, kol užpildysiu 124 milijonus vietų (panaudojimo rodiklis apie 0,96). Čia yra rezultatų lentelė, kiekviena eilutė yra CUDA branduolio iškvietimas, skirtas įterpti 4 milijonus naujų elementų į vieną maišos lentelę:
Naudojimo norma
Įterpimo trukmė 4 194 304 elementai
0,00
11,608448 ms (361,314798 mln. raktų/sek.)
0,03
11,751424 ms (356,918799 mln. raktų/sek.)
0,06
11,942592 ms (351,205515 mln. raktų/sek.)
0,09
12,081120 ms (347,178429 mln. raktų/sek.)
0,12
12,242560 ms (342,600233 mln. raktų/sek.)
0,16
12,396448 ms (338,347235 mln. raktų/sek.)
0,19
12,533024 ms (334,660176 mln. raktų/sek.)
0,22
12,703328 ms (330,173626 mln. raktų/sek.)
0,25
12,884512 ms (325,530693 mln. raktų/sek.)
0,28
13,033472 ms (321,810182 mln. raktų/sek.)
0,31
13,239296 ms (316,807174 mln. raktų/sek.)
0,34
13,392448 ms (313,184256 mln. raktų/sek.)
0,37
13,624000 ms (307,861434 mln. raktų/sek.)
0,41
13,875520 ms (302,280855 mln. raktų/sek.)
0,44
14,126528 ms (296,909756 mln. raktų/sek.)
0,47
14,399328 ms (291,284699 mln. raktų/sek.)
0,50
14,690304 ms (285,515123 mln. raktų/sek.)
0,53
15,039136 ms (278,892623 mln. raktų/sek.)
0,56
15,478656 ms (270,973402 mln. raktų/sek.)
0,59
15,985664 ms (262,379092 mln. raktų/sek.)
0,62
16,668673 ms (251,627968 mln. raktų/sek.)
0,66
17,587200 ms (238,486174 mln. raktų/sek.)
0,69
18,690048 ms (224,413765 mln. raktų/sek.)
0,72
20,278816 ms (206,831789 mln. raktų/sek.)
0,75
22,545408 ms (186,038058 mln. raktų/sek.)
0,78
26,053312 ms (160,989275 mln. raktų/sek.)
0,81
31,895008 ms (131,503463 mln. raktų/sek.)
0,84
42,103294 ms (99,619378 mln. raktų/sek.)
0,87
61,849056 ms (67,815164 mln. raktų/sek.)
0,90
105,695999 ms (39,682713 mln. raktų/sek.)
0,94
240,204636 ms (17,461378 mln. raktų/sek.)
Didėjant naudojimui, našumas mažėja. Daugeliu atvejų tai nepageidautina. Jei programa įterpia elementus į lentelę, o paskui juos atmeta (pavyzdžiui, skaičiuodama žodžius knygoje), tai nėra problema. Bet jei programa naudoja ilgalaikę maišos lentelę (pavyzdžiui, grafikos rengyklėje, kad saugotų netuščias vaizdų dalis, kuriose vartotojas dažnai įterpia ir ištrina informaciją), toks elgesys gali būti problemiškas.
Ir išmatavo maišos lentelės zondavimo gylį po 64 milijonų įterpimų (naudojimo koeficientas 0,5). Vidutinis gylis buvo 0,4774, todėl dauguma klavišų buvo geriausioje įmanomoje vietoje arba per vieną lizdą nuo geriausios padėties. Didžiausias zondavimo gylis buvo 60.
Tada išmatavau zondavimo gylį ant stalo su 124 milijonais įdėklų (naudojimo koeficientas 0,97). Vidutinis gylis jau buvo 10,1757, o didžiausias - 6474 (!!). Esant dideliam panaudojimo lygiui, tiesinio jutimo našumas žymiai sumažėja.
Geriausia, kad šios maišos lentelės naudojimo lygis būtų žemas. Bet tada padidiname našumą atminties suvartojimo sąskaita. Laimei, 32 bitų raktų ir reikšmių atveju tai gali būti pateisinama. Jei aukščiau pateiktame pavyzdyje lentelėje su 128 milijonais elementų išliksime 0,25 panaudojimo koeficientą, tada joje galime patalpinti ne daugiau kaip 32 milijonus elementų, o likę 96 milijonai lizdų bus prarasti - 8 baitai kiekvienai porai , 768 MB prarastos atminties.
Atkreipkite dėmesį, kad kalbame apie vaizdo plokštės atminties praradimą, kuri yra vertingesnis šaltinis nei sistemos atmintis. Nors dauguma šiuolaikinių stalinių kompiuterių grafikos plokščių, palaikančių CUDA, turi bent 4 GB atminties (rašymo metu NVIDIA 2080 Ti turi 11 GB), prarasti tokias sumas vis tiek nebūtų pats protingiausias sprendimas.
Vėliau parašysiu daugiau apie maišos lentelių kūrimą vaizdo plokštėms, kurios neturi problemų su zondavimo gyliu, taip pat būdus, kaip pakartotinai panaudoti ištrintus lizdus.
Garsavimo gylio matavimas
Norėdami nustatyti rakto tikrinimo gylį, galime išskirti rakto maišą (jo idealų lentelės indeksą) iš tikrojo lentelės indekso:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Dėl dvejetainių skaičių dvejetų papildymo magijos ir to, kad maišos lentelės talpa yra du laipsniai dviejų, šis metodas veiks net tada, kai pagrindinis indeksas perkeliamas į lentelės pradžią. Paimkime raktą, kurio maiša yra 1, bet yra įdėta į 3 lizdą. Tada lentelės talpa 4 gauname (3 — 1) & 3, kuris yra lygus 2.
išvada
Jei turite klausimų ar pastabų, rašykite man el arba atidarykite naują temą.
Šis kodas buvo parašytas įkvėpus puikių straipsnių:
Ateityje ir toliau rašysiu apie maišos lentelės diegimus vaizdo plokštėms ir analizuosiu jų veikimą. Mano planuose yra grandinės, Robino Hudo maišos ir gegutės maišos naudojimas naudojant atomines operacijas duomenų struktūrose, kurios yra draugiškos GPU.
Šaltinis: www.habr.com