Github-en argitaratu nuen .
GPU hash taula soil bat da, segundoko ehunka milioi txertatze prozesatzeko gai dena. Nire NVIDIA GTX 1060 ordenagailu eramangarrian, kodeak ausaz sortutako 64 milioi gako-balio bikote txertatzen ditu 210 ms inguru eta 32 milioi bikote kentzen ditu 64 ms inguru.
Hau da, ordenagailu eramangarri baten abiadura gutxi gorabehera 300 milioi txertatze/seg eta 500 milioi ezabatze/seg da.
Taula CUDAn idatzita dago, nahiz eta teknika bera HLSL edo GLSLra aplika daitekeen. Inplementazioak hainbat muga ditu bideo-txartel batean errendimendu handia bermatzeko:
- 32 biteko gakoak eta balio berdinak soilik prozesatzen dira.
- Hash taulak tamaina finkoa du.
- Eta tamaina horrek potentziaren bi berdina izan behar du.
Gako eta balioetarako, mugatzaile-markatzaile soil bat erreserbatu behar duzu (goiko kodean hau 0xffffffff da).
Sarrailarik gabeko hash taula
Hash taulak helbide irekia erabiltzen du , hau da, memorian gordetzen den eta cache-ren errendimendu handiagoa duen gako-balio bikoteen array bat besterik ez da. Ezin da gauza bera esan kateatzeari dagokionez, hau da, erakuslea loturiko zerrenda batean bilatzea dakar. Hash taula elementuak biltegiratzeko array sinple bat da KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Taularen tamaina biko potentzia da, ez zenbaki lehen bat, instrukzio azkar bat nahikoa baita pow2/AND maskara aplikatzeko, baina moduluaren operadorea askoz motelagoa da. Hau garrantzitsua da zundaketa linealaren kasuan, izan ere, taula linealeko bilaketa batean zirrikitu indizea zirrikitu bakoitzean bildu behar da. Eta ondorioz, eragiketaren kostua modulu gehitzen da zirrikitu bakoitzean.
Taulak elementu bakoitzaren gakoa eta balioa soilik gordetzen ditu, ez gakoaren hash bat. Taulak 32 biteko gakoak soilik gordetzen dituenez, hash-a oso azkar kalkulatzen da. Goiko kodeak Murmur3 hash-a erabiltzen du, txandaka, XOR eta biderketa batzuk bakarrik egiten dituena.
Hash-taulak blokeatzeko babes-teknikak erabiltzen ditu, memoria-ordenatik independenteak direnak. Idazketa-eragiketa batzuek beste eragiketa batzuen ordena apurtzen badute ere, hash taulak egoera egokia mantenduko du. Honetaz hitz egingo dugu jarraian. Teknikak bikain funtzionatzen du milaka hari aldi berean exekutatzen dituzten bideo-txartelekin.
Hash taulako gakoak eta balioak hasieratzen dira hutsik egoteko.
Kodea alda daiteke 64 biteko gakoak eta balioak kudeatzeko. Gakoek irakurketa, idazketa eta konparatzeko eta trukatzeko eragiketa atomikoak behar dituzte. Eta balioek irakurketa eta idazketa atomikoak behar dituzte. Zorionez, CUDAn, 32 eta 64 biteko balioetarako irakurketa-idazketa eragiketak atomikoak dira, naturalki lerrokatuta dauden bitartean (ikus behean). ), eta bideo-txartel modernoek 64 biteko konparatzeko eta trukatzeko eragiketa atomikoak onartzen dituzte. Jakina, 64 bitera pasatzean, errendimendua pixka bat jaitsiko da.
Hash taularen egoera
Hash taula bateko gako-balio bikote bakoitzak lau egoera hauetako bat izan dezake:
- Gakoa eta balioa hutsik daude. Egoera honetan, hash taula hasieratzen da.
- Gakoa idatzita dago, baina balioa oraindik ez da idatzi. Momentu honetan beste hari bat datuak irakurtzen ari bada, hutsik itzultzen da. Hau normala da, gauza bera gertatuko zen beste exekuzio-hari batek apur bat lehenago funtzionatu izan balu, eta aldibereko datu-egituraz ari gara.
- Gakoa eta balioa erregistratzen dira.
- Balioa beste exekuzio hari batzuentzat eskuragarri dago, baina gakoa ez dago oraindik. Hau gerta daiteke CUDA programazio-ereduak ordenatutako memoria-eredu baxua duelako. Hau normala da; edonola ere, gakoa hutsik dago oraindik, nahiz eta balioa jada ez den.
Γabardura garrantzitsu bat da gakoa zirrikituan idatzitakoan, jada ez dela mugitzen; gakoa ezabatu bada ere, honetaz hitz egingo dugu jarraian.
Hash-taularen kodeak ordena baxuko memoria-ereduekin ere funtzionatzen du, zeinetan memoria irakurtzeko eta idazteko ordena ezezaguna den. Hash-taula batean txertatzea, bilatu eta ezabatzeari begiratzen diogunean, gogoratu gako-balio bikote bakoitza goian deskribatutako lau egoeretako batean dagoela.
Hash taula batean txertatzea
Hash taula batean gako-balio bikoteak txertatzen dituen CUDA funtzioak itxura hau du:
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);
}
}
Gako bat txertatzeko, kodea hash-taularen arrayaren bidez errepikatzen da, txertatutako gakoaren hashetik hasita. Array-ko zirrikitu bakoitzak konparatzeko eta trukatzeko eragiketa atomiko bat egiten du, zirrikitu horretako giltza hutsarekin alderatzen duena. Bat-etortze bat antzematen bada, zirrikituan dagoen gakoa txertatutako giltzarekin eguneratzen da eta, ondoren, jatorrizko zirrikituaren gakoa itzultzen da. Jatorrizko gako hau hutsik bazegoen edo sartutako gakoarekin bat zetorren, orduan kodeak txertatzeko zirrikitu egoki bat aurkitu zuen eta txertatutako balioa zirrikituan sartu zuen.
Kernel-dei batean bada gpu_hashtable_insert() gako berdina duten hainbat elementu daude, gero haien balioetako edozein gakoen zirrikituan idatz daiteke. Hau normaltzat jotzen da: deian zehar gako-balioen idazketa bat arrakastatsua izango da, baina hori guztia exekuzioaren hainbat haritan paraleloan gertatzen denez, ezin dugu aurreikusi zein memoriaren idazketa izango den azkena.
Hash taularen bilaketa
Gakoak bilatzeko kodea:
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);
}
}
Taula batean gordetako gako baten balioa aurkitzeko, matrizean zehar iteratuko dugu bilatzen ari garen gakoaren hashetik hasita. Zirrikitu bakoitzean, gakoa bilatzen ari garen gakoa den egiaztatzen dugu, eta hala bada, bere balioa itzuliko dugu. Gakoa hutsik dagoen ere egiaztatzen dugu, eta hala bada, bilaketa bertan behera uzten dugu.
Ezin badugu gakoa aurkitzen, kodeak balio huts bat itzultzen du.
Bilaketa-eragiketa hauek guztiak aldi berean egin daitezke txertatzeen eta ezabaketen bidez. Taulan bikote bakoitzak goian deskribatutako lau egoeretako bat izango du fluxurako.
Hash taula batean ezabatzen
Gakoak ezabatzeko kodea:
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);
}
}
Gako bat ezabatzea ezohiko modu batean egiten da: gakoa taulan uzten dugu eta bere balioa (ez gakoa bera) hutsik markatuko dugu. Kode hau oso antzekoa da lookup(), gako batean bat-etortze bat aurkitzen denean bere balioa hutsik egiten duela izan ezik.
Goian esan bezala, gako bat zirrikitu batean idatzitakoan, jada ez da mugitzen. Elementu bat taulatik ezabatzen denean ere, gakoak bere lekuan jarraitzen du, bere balioa hutsik geratzen da. Horrek esan nahi du ez dugula idazketa atomikorik erabili behar zirrikituaren baliorako, ez baitu axola uneko balioa hutsik dagoen ala ez - hutsik geratuko da oraindik.
Hash taula baten tamaina aldatu
Hash-taula baten tamaina alda dezakezu taula handiago bat sortuz eta taula zaharreko elementu hutsak txertatuz. Ez nuen funtzionalitate hau inplementatu lagin kodea sinplea mantendu nahi nuelako. Gainera, CUDA programetan, memoria esleipena ostalariaren kodean egiten da sarritan CUDA nukleoan baino.
Artikulua Blokeo bidez babestutako datu-egitura hori nola aldatu deskribatzen du.
Lehiakortasuna
Goiko funtzio-kode zatietan gpu_hashtable_insert(), _lookup() ΠΈ _delete() gako-balio bikote bat prozesatu aldi berean. Eta baxuago gpu_hashtable_insert(), _lookup() ΠΈ _delete() pare-matrize bat prozesatu paraleloan, bikote bakoitza GPU exekuzio-hari bereizi batean:
// 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);
}
}
Blokeo-erresistentzia hash taulak aldibereko txertaketak, bilaketak eta ezabaketak onartzen ditu. Gako-balio bikoteak beti lau egoeratako batean daudenez eta gakoak ez direnez mugitzen, taulak zuzentasuna bermatzen du eragiketa mota desberdinak aldi berean erabiltzen badira ere.
Hala ere, txertatze eta ezabatze sorta bat paraleloan prozesatzen badugu eta bikoteen sarrerako matrizeak gako bikoiztuak baditu, ezin izango dugu aurreikusi zein bikote "irabaziko" duten; hash taulan idatziko da azkena. Demagun txertatzeko kodea deitzen diogula bikoteen sarrerako matrize batekin A/0 B/1 A/2 C/3 A/4. Kodea osatzen denean, bikoteka B/1 ΠΈ C/3 taulan egotea bermatuta dago, baina aldi berean bikoteren bat agertuko da bertan A/0, A/2 edo A/4. Arazo bat izan daiteke edo ez; dena aplikazioaren araberakoa da. Agian aldez aurretik jakin dezakezu sarrerako matrizean bikoiztutako gakorik ez dagoela, edo agian ez zaizu axola zein balio idatzi den azkena.
Zuretzat arazoa bada, bikoiztutako bikoteak CUDA sistema dei desberdinetan banatu behar dituzu. CUDAn, nukleoa deitzen duen edozein eragiketa hurrengo nukleoaren deia baino lehen amaitzen da beti (gutxienez hari baten barruan. Hari ezberdinetan, nukleoak paraleloan exekutatzen dira). Goiko adibidean, nukleo bati deitzen badiozu A/0 B/1 A/2 C/3, eta bestearekin A/4, gero giltza A balioa jasoko du 4.
Orain hitz egin dezagun funtzioek behar duten ala ez lookup() ΠΈ delete() Erabili erakusle arrunta edo lurrunkorra hash taulako bikote-matrize baterako. Hau dio:
Konpilatzaileak memoria global edo partekatuan irakurketak eta idazketak optimizatzea aukera dezake... Optimizazio hauek desgaitu daitezke gako-hitza erabiliz.
volatile: ... aldagai honi buruzko edozein erreferentzia memoria erreal batean irakurtzeko edo idazteko instrukzio batean biltzen da.
Zuzentasun kontuek ez dute aplikaziorik behar volatile. Exekuzio-hariak aurreko irakurketa-eragiketa bateko cache-ko balio bat erabiltzen badu, apur bat zaharkitutako informazioa erabiliko du. Hala ere, hau hash taularen egoera zuzenaren informazioa da nukleoaren deiaren une jakin batean. Azken informazioa erabili behar baduzu, aurkibidea erabil dezakezu volatile, baina gero errendimendua apur bat jaitsiko da: nire proben arabera, 32 milioi elementu ezabatzean, abiadura 500 milioi ezabatze/seg izatetik 450 milioi ezabatze/seg izatera pasa zen.
produktibitatea
64 milioi elementu sartu eta horietako 32 milioi ezabatzeko proban, arteko lehia std::unordered_map eta ez dago ia hash taularik GPUrako:
std::unordered_map 70 ms eman ditu elementuak sartzen eta kentzen eta gero askatzen unordered_map (Milioika elementu kentzeak denbora asko behar du, barrutik unordered_map memoria esleipen anitz egiten dira). Egia esanda, std:unordered_map murrizketa guztiz desberdinak. PUZaren exekuzio hari bakarra da, edozein tamainatako gako-balioak onartzen ditu, erabilera-tasa altuetan ondo funtzionatzen du eta errendimendu egonkorra erakusten du hainbat ezabapenen ondoren.
GPUrako eta programen arteko komunikaziorako hash taularen iraupena 984 ms izan zen. Taula memorian jartzen eta ezabatzen igarotako denbora barne hartzen du (1 GB memoria behin esleituz, denbora pixka bat hartzen duena CUDAn), elementuak txertatu eta ezabatuz eta horien gainean errepikatuz. Bideo-txarteleko memoriara egindako kopia guztiak ere kontuan hartzen dira.
Hash taulak 271 ms behar izan zituen osatzeko. Horrek bideo-txartelak elementuak txertatzen eta ezabatzen igarotzen duen denbora barne hartzen du, eta ez du kontuan hartzen memorian kopiatzen eta emaitza taularen gainean errepikatzen emandako denbora. GPU taula denbora luzez bizi bada, edo hash taula bideo-txartelaren memorian osorik badago (adibidez, beste GPU kodeak eta ez prozesadore zentralak erabiliko duen hash taula bat sortzeko), orduan probaren emaitza garrantzitsua da.
Bideo-txartel baten hash-taulak errendimendu handia erakusten du errendimendu handiko eta paralelizazio aktiboaren ondorioz.
Mugak
Hash-taularen arkitekturak arazo batzuk ditu kontutan izan:
- Zundaketa lineala multzokatzeak oztopatzen du, eta horrek taulako giltzak ezin hobeto kokatzea eragiten du.
- Teklak ez dira kentzen funtzioa erabiliz
deleteeta denborarekin mahaia nahasten dute.
Ondorioz, hash-taula baten errendimendua pixkanaka honda daiteke, batez ere denbora luzez existitzen bada eta txertatze eta ezabatze ugari baditu. Desabantaila hauek arintzeko modu bat erabilera-tasa nahiko baxua duen taula berri batean birpasatzea da eta kendutako gakoak birpasatzea zehar iragaztea da.
Azaldutako gaiak ilustratzeko, goiko kodea erabiliko dut 128 milioi elementu dituen taula bat sortzeko eta 4 milioi elementu igarotzeko 124 milioi zirrikitu arte (0,96 inguruko erabilera-tasa) bete arte. Hona hemen emaitza-taula, errenkada bakoitza CUDA kernel-dei bat da 4 milioi elementu berri hash taula batean txertatzeko:
Erabilera tasa
Txertazioaren iraupena 4 elementu
0,00
11,608448 ms (361,314798 milioi gako/seg.)
0,03
11,751424 ms (356,918799 milioi gako/seg.)
0,06
11,942592 ms (351,205515 milioi gako/seg.)
0,09
12,081120 ms (347,178429 milioi gako/seg.)
0,12
12,242560 ms (342,600233 milioi gako/seg.)
0,16
12,396448 ms (338,347235 milioi gako/seg.)
0,19
12,533024 ms (334,660176 milioi gako/seg.)
0,22
12,703328 ms (330,173626 milioi gako/seg.)
0,25
12,884512 ms (325,530693 milioi gako/seg.)
0,28
13,033472 ms (321,810182 milioi gako/seg.)
0,31
13,239296 ms (316,807174 milioi gako/seg.)
0,34
13,392448 ms (313,184256 milioi gako/seg.)
0,37
13,624000 ms (307,861434 milioi gako/seg.)
0,41
13,875520 ms (302,280855 milioi gako/seg.)
0,44
14,126528 ms (296,909756 milioi gako/seg.)
0,47
14,399328 ms (291,284699 milioi gako/seg.)
0,50
14,690304 ms (285,515123 milioi gako/seg.)
0,53
15,039136 ms (278,892623 milioi gako/seg.)
0,56
15,478656 ms (270,973402 milioi gako/seg.)
0,59
15,985664 ms (262,379092 milioi gako/seg.)
0,62
16,668673 ms (251,627968 milioi gako/seg.)
0,66
17,587200 ms (238,486174 milioi gako/seg.)
0,69
18,690048 ms (224,413765 milioi gako/seg.)
0,72
20,278816 ms (206,831789 milioi gako/seg.)
0,75
22,545408 ms (186,038058 milioi gako/seg.)
0,78
26,053312 ms (160,989275 milioi gako/seg.)
0,81
31,895008 ms (131,503463 milioi gako/seg.)
0,84
42,103294 ms (99,619378 milioi gako/seg.)
0,87
61,849056 ms (67,815164 milioi gako/seg.)
0,90
105,695999 ms (39,682713 milioi gako/seg.)
0,94
240,204636 ms (17,461378 milioi gako/seg.)
Erabilera handitu ahala, errendimendua gutxitzen da. Hau ez da desiragarria kasu gehienetan. Aplikazio batek elementuak taula batean txertatzen baditu eta gero baztertzen baditu (adibidez, liburu bateko hitzak zenbatzean), ez da arazoa. Baina aplikazioak iraupen luzeko hash-taula erabiltzen badu (adibidez, editore grafiko batean hutsik gabeko irudien zatiak gordetzeko, non erabiltzaileak maiz txertatzen eta ezabatzen du informazioa), orduan jokabide hori arazotsua izan daiteke.
Eta hash-taularen zundaketaren sakonera neurtu zuen 64 milioi txertatzeren ondoren (erabilpen-faktorea 0,5). Batez besteko sakonera 0,4774koa zen, beraz, giltza gehienak ahalik eta zirrikitu onenean zeuden edo posizio onenetik urrun zeuden zirrikitua. Soinu-sakonera maximoa 60 izan zen.
Ondoren, zundaketa-sakonera 124 milioi txertatze dituen mahai batean neurtu nuen (erabilera-faktorea 0,97). Batez besteko sakonera jada 10,1757koa zen, eta maximoa - 6474 (!!). Sentsazio linealaren errendimendua nabarmen jaisten da erabilera-tasa altuetan.
Hobe da hash-taularen erabilera-tasa baxua izatea. Baina gero errendimendua handitzen dugu memoria kontsumoaren kontura. Zorionez, 32 biteko gako eta balioen kasuan, hori justifika daiteke. Goiko adibidean, 128 milioi elementu dituen taula batean, 0,25eko erabilera-faktorea mantentzen badugu, orduan ezin ditugu 32 milioi elementu baino gehiago jarri bertan, eta gainerako 96 milioi zirrikituak galduko dira - 8 byte bikote bakoitzeko. , 768 MB galdutako memoria.
Kontuan izan bideo-txartelen memoria galtzeaz ari garela, sistemaren memoria baino baliabide baliotsuagoa baita. CUDA onartzen duten mahaigaineko txartel grafiko moderno gehienek gutxienez 4 GB memoria badute ere (idazteko unean, NVIDIA 2080 Ti-k 11 GB ditu), oraindik ez litzateke erabakirik zuhurrena izango kopuru horiek galtzea.
Geroago gehiago idatziko dut zundaketaren sakonerarekin arazorik ez duten bideo-txartelen hash taulak sortzeari buruz, baita ezabatutako zirrikituak berrerabiltzeko moduei buruz ere.
Soinuaren sakoneraren neurketa
Gako baten zundaketa-sakonera zehazteko, gakoaren hash (bere taula-indize ideala) atera dezakegu bere benetako taula-indizetik:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Biren bi osagarrien zenbaki bitarren magia dela eta eta hash-taularen ahalmena biren potentziarako bikoa denez, ikuspegi honek funtzionatuko du gako-indizea taularen hasierara eraman arren. Har dezagun 1era hashatutako gako bat, baina 3. zirrikituan txertatzen den. Ondoren, 4. edukiera duen taula baterako, lortuko dugu (3 β 1) & 3, 2-ren baliokidea dena.
Ondorioa
Galderak edo iruzkinak badituzu, mesedez, bidali mezu elektronikoa helbide honetara edo ireki gai berri bat.
Kode hau artikulu bikainetan inspiratuta idatzi zen:
Etorkizunean, bideo-txarteletarako hash taula inplementazioei buruz idazten eta haien errendimendua aztertzen jarraituko dut. Nire planak kateatzea, Robin Hood hashinga eta kukua hashinga dira GPU errespetatzen duten datu egituretan eragiketa atomikoak erabiliz.
Iturria: www.habr.com