E postova në Github .
Është një tabelë e thjeshtë hash GPU e aftë për të përpunuar qindra miliona inserte në sekondë. Në laptopin tim NVIDIA GTX 1060, kodi fut 64 milionë çifte të vlerave kyçe të krijuara rastësisht në rreth 210 ms dhe heq 32 milionë çifte në rreth 64 ms.
Kjo do të thotë, shpejtësia në një laptop është afërsisht 300 milionë inserte/sek dhe 500 milionë fshirje/sek.
Tabela është shkruar në CUDA, megjithëse e njëjta teknikë mund të zbatohet për HLSL ose GLSL. Zbatimi ka disa kufizime për të siguruar performancë të lartë në një kartë video:
- Përpunohen vetëm çelësat 32-bit dhe të njëjtat vlera.
- Tabela hash ka një madhësi fikse.
- Dhe kjo madhësi duhet të jetë e barabartë me dy me fuqinë.
Për çelësat dhe vlerat, duhet të rezervoni një shënues të thjeshtë ndarës (në kodin e mësipërm është 0xffffffff).
Tavolinë hash pa bravë
Tabela hash përdor adresimin e hapur me , domethënë, është thjesht një grup çiftesh vlerash çelësi që ruhet në memorie dhe ka performancë superiore të cache-it. Nuk mund të thuhet e njëjta gjë për zinxhirin, i cili përfshin kërkimin e një treguesi në një listë të lidhur. Një tabelë hash është një grup i thjeshtë që ruan elementë KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Madhësia e tabelës është fuqia e dy, jo një numër i thjeshtë, sepse mjafton një instruksion i shpejtë për të aplikuar maskën pow2/AND, por operatori i modulit është shumë më i ngadalshëm. Kjo është e rëndësishme në rastin e kërkimit linear, pasi në një kërkim linear tabele, indeksi i slotit duhet të mbështillet në secilën slot. Dhe si rezultat, kostoja e operacionit shtohet modul në çdo slot.
Tabela ruan vetëm çelësin dhe vlerën për çdo element, jo një hash të çelësit. Meqenëse tabela ruan vetëm çelësat 32-bit, hash-i llogaritet shumë shpejt. Kodi i mësipërm përdor hash-in Murmur3, i cili kryen vetëm disa ndërrime, XOR dhe shumëzime.
Tabela hash përdor teknika të mbrojtjes së bllokimit që janë të pavarura nga renditja e kujtesës. Edhe nëse disa operacione shkrimi prishin rendin e operacioneve të tjera të tilla, tabela hash do të vazhdojë të ruajë gjendjen e duhur. Ne do të flasim për këtë më poshtë. Teknika funksionon shkëlqyeshëm me kartat video që drejtojnë mijëra tema njëkohësisht.
Çelësat dhe vlerat në tabelën hash janë inicializuar për të zbrazur.
Kodi mund të modifikohet për të trajtuar gjithashtu çelësat dhe vlerat 64-bit. Çelësat kërkojnë operacione atomike leximi, shkrimi dhe krahasimi dhe shkëmbimi. Dhe vlerat kërkojnë operacione atomike të leximit dhe shkrimit. Për fat të mirë, në CUDA, operacionet e leximit-shkrimit për vlerat 32- dhe 64-bit janë atomike për sa kohë që ato janë të lidhura natyrshëm (shih më poshtë). ), dhe kartat moderne video mbështesin operacionet e krahasimit dhe shkëmbimit atomik 64-bit. Sigurisht, kur kaloni në 64 bit, performanca do të ulet pak.
Gjendja e tabelës hash
Çdo çift çelës-vlerë në një tabelë hash mund të ketë një nga katër gjendjet:
- Çelësi dhe vlera janë bosh. Në këtë gjendje, tabela hash inicializohet.
- Çelësi është shkruar, por vlera ende nuk është shkruar. Nëse një thread tjetër aktualisht po lexon të dhëna, atëherë ai kthehet bosh. Kjo është normale, e njëjta gjë do të kishte ndodhur nëse një fije tjetër ekzekutimi do të kishte funksionuar pak më herët, dhe ne po flasim për një strukturë të njëkohshme të të dhënave.
- Të dyja, çelësi dhe vlera regjistrohen.
- Vlera është e disponueshme për fijet e tjera të ekzekutimit, por çelësi nuk është ende. Kjo mund të ndodhë sepse modeli i programimit CUDA ka një model memorie të renditur lirshëm. Kjo është normale; në çdo rast, çelësi është ende bosh, edhe nëse vlera nuk është më e tillë.
Një nuancë e rëndësishme është që pasi çelësi të jetë shkruar në slot, ai nuk lëviz më - edhe nëse çelësi fshihet, ne do të flasim për këtë më poshtë.
Kodi i tabelës hash funksionon edhe me modele memorie të renditura lirshëm në të cilat rendi në të cilin lexohet dhe shkruhet memoria është i panjohur. Ndërsa shikojmë futjen, kërkimin dhe fshirjen në një tabelë hash, mbani mend se çdo çift çelës-vlerë është në një nga katër gjendjet e përshkruara më sipër.
Futja në një tabelë hash
Funksioni CUDA që fut çiftet çelës-vlerë në një tabelë hash duket si ky:
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);
}
}
Për të futur një çelës, kodi përsëritet përmes grupit të tabelës hash duke filluar me hash-in e çelësit të futur. Çdo slot në grup kryen një operacion atomik krahasim-dhe-këmbim që krahason çelësin në atë slot me zbrazjen. Nëse zbulohet një mospërputhje, çelësi në slot përditësohet me çelësin e futur dhe më pas kthehet çelësi origjinal i slotit. Nëse ky çelës origjinal ishte bosh ose përputhej me çelësin e futur, atëherë kodi gjeti një vend të përshtatshëm për futje dhe futi vlerën e futur në slot.
Nëse në një thirrje kernel gpu_hashtable_insert() ka shumë elementë me të njëjtin çelës, atëherë çdo vlerë e tyre mund të shkruhet në folenë e çelësit. Kjo konsiderohet normale: një nga shkrimet me vlerë kyçe gjatë thirrjes do të ketë sukses, por meqenëse e gjithë kjo ndodh paralelisht brenda disa fijeve të ekzekutimit, ne nuk mund të parashikojmë se cili shkrim memorie do të jetë i fundit.
Kërkimi i tabelës hash
Kodi për kërkimin e çelësave:
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);
}
}
Për të gjetur vlerën e një çelësi të ruajtur në një tabelë, ne përsërisim përmes grupit duke filluar me hash-in e çelësit që kërkojmë. Në çdo slot, kontrollojmë nëse çelësi është ai që kërkojmë dhe nëse po, ia kthejmë vlerën. Ne gjithashtu kontrollojmë nëse çelësi është bosh dhe nëse po, ne e anulojmë kërkimin.
Nëse nuk mund ta gjejmë çelësin, kodi kthen një vlerë boshe.
Të gjitha këto operacione kërkimi mund të kryhen njëkohësisht përmes futjeve dhe fshirjeve. Çdo çift në tabelë do të ketë një nga katër gjendjet e përshkruara më sipër për rrjedhën.
Fshirja në një tabelë hash
Kodi për fshirjen e çelësave:
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);
}
}
Fshirja e një çelësi bëhet në një mënyrë të pazakontë: e lëmë çelësin në tabelë dhe shënojmë vlerën e tij (jo vetë çelësin) si bosh. Ky kod është shumë i ngjashëm me lookup(), përveç se kur një ndeshje gjendet në një çelës, ai e bën vlerën e tij boshe.
Siç u përmend më lart, pasi një çelës shkruhet në një vend, ai nuk zhvendoset më. Edhe kur një element fshihet nga tabela, çelësi mbetet në vend, vlera e tij thjesht bëhet bosh. Kjo do të thotë që ne nuk kemi nevojë të përdorim një operacion shkrimi atomik për vlerën e slotit, sepse nuk ka rëndësi nëse vlera aktuale është bosh apo jo - ajo do të mbetet ende bosh.
Ndryshimi i madhësisë së një tabele hash
Ju mund të ndryshoni madhësinë e një tabele hash duke krijuar një tabelë më të madhe dhe duke futur elementë jo bosh nga tabela e vjetër në të. Unë nuk e zbatova këtë funksionalitet sepse doja ta mbaja të thjeshtë kodin e mostrës. Për më tepër, në programet CUDA, shpërndarja e memories shpesh bëhet në kodin pritës dhe jo në kernelin CUDA.
Në artikull përshkruan se si të modifikohet një strukturë e tillë e të dhënave e mbrojtur nga bllokimi.
Konkurrueshmëria
Në copat e kodit të funksionit të mësipërm gpu_hashtable_insert(), _lookup() и _delete() përpunoni një çift çelës-vlerë në të njëjtën kohë. Dhe më e ulët gpu_hashtable_insert(), _lookup() и _delete() përpunoni një grup çiftesh paralelisht, secili çift në një fije të veçantë ekzekutimi GPU:
// 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);
}
}
Tabela hash rezistente ndaj bllokimit mbështet futjet, kërkimet dhe fshirjet e njëkohshme. Për shkak se çiftet çelës-vlerë janë gjithmonë në një nga katër gjendjet dhe çelësat nuk lëvizin, tabela garanton korrektësi edhe kur lloje të ndryshme operacionesh përdoren njëkohësisht.
Megjithatë, nëse përpunojmë një grup futjesh dhe fshirjesh paralelisht, dhe nëse grupi hyrës i çifteve përmban çelësa dublikatë, atëherë nuk do të jemi në gjendje të parashikojmë se cilat çifte do të "fitojnë" - do të shkruhet në tabelën hash të fundit. Le të themi se kemi thirrur kodin e futjes me një grup hyrës çiftesh A/0 B/1 A/2 C/3 A/4. Kur kodi përfundon, çiftohet B/1 и C/3 janë të garantuara të jenë të pranishme në tabelë, por në të njëjtën kohë ndonjë nga çiftet do të shfaqet në të A/0, A/2 ose A/4. Ky mund të jetë ose jo problem - gjithçka varet nga aplikacioni. Ju mund ta dini paraprakisht se nuk ka çelësa dublikatë në grupin e hyrjes, ose mund të mos ju interesojë se cila vlerë është shkruar e fundit.
Nëse ky është një problem për ju, atëherë duhet të ndani çiftet e kopjuara në thirrje të ndryshme të sistemit CUDA. Në CUDA, çdo operacion që thërret kernelin përfundon gjithmonë përpara thirrjes tjetër të kernelit (të paktën brenda një thread. Në thread të ndryshëm, kernelet ekzekutohen paralelisht). Në shembullin e mësipërm, nëse thërrisni një kernel me A/0 B/1 A/2 C/3, dhe tjetra me A/4, pastaj çelësi A do të marrë vlerën 4.
Tani le të flasim nëse funksionet duhet lookup() и delete() përdorni një tregues të thjeshtë ose të paqëndrueshëm për një grup çiftesh në tabelën hash. shprehet se:
Përpiluesi mund të zgjedhë të optimizojë leximet dhe shkrimet në kujtesën globale ose të përbashkët... Këto optimizime mund të çaktivizohen duke përdorur fjalën kyçe
volatile: ... çdo referencë për këtë variabël përpilohet në një instruksion leximi ose shkrimi në memorie reale.
Konsideratat e korrektësisë nuk kërkojnë aplikim volatile. Nëse filli i ekzekutimit përdor një vlerë të memorizuar nga një operacion i mëparshëm leximi, atëherë do të përdorë informacion paksa të vjetëruar. Por megjithatë, ky është informacion nga gjendja e saktë e tabelës hash në një moment të caktuar të thirrjes së kernelit. Nëse keni nevojë të përdorni informacionin më të fundit, mund të përdorni indeksin volatile, por më pas performanca do të ulet paksa: sipas testeve të mia, kur fshini 32 milionë elementë, shpejtësia u ul nga 500 milionë fshirje/sek në 450 milionë fshirje/sek.
prodhimtari
Në testin për futjen e 64 milionë elementëve dhe fshirjen e 32 milionë prej tyre, konkurrenca mes std::unordered_map dhe praktikisht nuk ka asnjë tabelë hash për GPU:
std::unordered_map shpenzoi 70 ms duke futur dhe hequr elementë dhe më pas duke i çliruar ato unordered_map (Heqja e miliona elementeve kërkon shumë kohë, sepse brenda unordered_map bëhen alokime të shumta memorie). Me thënë të drejtën, std:unordered_map kufizime krejtësisht të ndryshme. Është një fije e vetme e ekzekutimit të CPU-së, mbështet vlerat kyçe të çdo madhësie, performon mirë me shkallë të lartë përdorimi dhe tregon performancë të qëndrueshme pas fshirjeve të shumta.
Kohëzgjatja e tabelës hash për GPU-në dhe komunikimin ndër-program ishte 984 ms. Kjo përfshin kohën e kaluar për vendosjen e tabelës në memorie dhe fshirjen e saj (ndarja e 1 GB memorie një herë, e cila kërkon pak kohë në CUDA), futjen dhe fshirjen e elementeve dhe përsëritjen mbi to. Të gjitha kopjet në dhe nga memoria e kartës video merren gjithashtu parasysh.
Vetë tabela hash mori 271 ms për t'u plotësuar. Kjo përfshin kohën e kaluar nga karta video për futjen dhe fshirjen e elementeve, dhe nuk merr parasysh kohën e kaluar për kopjimin në memorie dhe përsëritjen mbi tabelën që rezulton. Nëse tabela GPU jeton për një kohë të gjatë, ose nëse tabela hash gjendet tërësisht në kujtesën e kartës video (për shembull, për të krijuar një tabelë hash që do të përdoret nga kodi tjetër GPU dhe jo nga procesori qendror), atëherë rezultati i testit është i rëndësishëm.
Tabela hash për një kartë video tregon performancë të lartë për shkak të xhiros së lartë dhe paralelizimit aktiv.
Kufizimet
Arkitektura e tabelës hash ka disa çështje që duhen pasur parasysh:
- Kërkimi linear pengohet nga grumbullimi, gjë që bën që çelësat në tabelë të vendosen më pak se në mënyrë perfekte.
- Tastet nuk hiqen duke përdorur funksionin
deletedhe me kalimin e kohës rrëmojnë tryezën.
Si rezultat, performanca e një tabele hash mund të degradojë gradualisht, veçanërisht nëse ekziston për një kohë të gjatë dhe ka inserte dhe fshirje të shumta. Një mënyrë për të zbutur këto disavantazhe është rihapja në një tabelë të re me një shkallë përdorimi mjaft të ulët dhe filtrimi i çelësave të hequr gjatë rihapjes.
Për të ilustruar çështjet e përshkruara, do të përdor kodin e mësipërm për të krijuar një tabelë me 128 milionë elementë dhe do të kaloj nëpër 4 milionë elementë derisa të plotësoj 124 milionë slota (shkalla e përdorimit prej rreth 0,96). Këtu është tabela e rezultateve, çdo rresht është një thirrje e kernelit CUDA për të futur 4 milionë elementë të rinj në një tabelë hash:
Shkalla e përdorimit
Kohëzgjatja e futjes 4 elemente
0,00
11,608448 ms (361,314798 milion çelësa/sek.)
0,03
11,751424 ms (356,918799 milion çelësa/sek.)
0,06
11,942592 ms (351,205515 milion çelësa/sek.)
0,09
12,081120 ms (347,178429 milion çelësa/sek.)
0,12
12,242560 ms (342,600233 milion çelësa/sek.)
0,16
12,396448 ms (338,347235 milion çelësa/sek.)
0,19
12,533024 ms (334,660176 milion çelësa/sek.)
0,22
12,703328 ms (330,173626 milion çelësa/sek.)
0,25
12,884512 ms (325,530693 milion çelësa/sek.)
0,28
13,033472 ms (321,810182 milion çelësa/sek.)
0,31
13,239296 ms (316,807174 milion çelësa/sek.)
0,34
13,392448 ms (313,184256 milion çelësa/sek.)
0,37
13,624000 ms (307,861434 milion çelësa/sek.)
0,41
13,875520 ms (302,280855 milion çelësa/sek.)
0,44
14,126528 ms (296,909756 milion çelësa/sek.)
0,47
14,399328 ms (291,284699 milion çelësa/sek.)
0,50
14,690304 ms (285,515123 milion çelësa/sek.)
0,53
15,039136 ms (278,892623 milion çelësa/sek.)
0,56
15,478656 ms (270,973402 milion çelësa/sek.)
0,59
15,985664 ms (262,379092 milion çelësa/sek.)
0,62
16,668673 ms (251,627968 milion çelësa/sek.)
0,66
17,587200 ms (238,486174 milion çelësa/sek.)
0,69
18,690048 ms (224,413765 milion çelësa/sek.)
0,72
20,278816 ms (206,831789 milion çelësa/sek.)
0,75
22,545408 ms (186,038058 milion çelësa/sek.)
0,78
26,053312 ms (160,989275 milion çelësa/sek.)
0,81
31,895008 ms (131,503463 milion çelësa/sek.)
0,84
42,103294 ms (99,619378 milion çelësa/sek.)
0,87
61,849056 ms (67,815164 milion çelësa/sek.)
0,90
105,695999 ms (39,682713 milion çelësa/sek.)
0,94
240,204636 ms (17,461378 milion çelësa/sek.)
Me rritjen e përdorimit, performanca zvogëlohet. Kjo nuk është e dëshirueshme në shumicën e rasteve. Nëse një aplikacion fut elementë në një tabelë dhe më pas i hedh ato (për shembull, kur numëron fjalët në një libër), atëherë ky nuk është problem. Por nëse aplikacioni përdor një tabelë hash jetëgjatë (për shembull, në një redaktues grafik për të ruajtur pjesët jo bosh të imazheve ku përdoruesi shpesh fut dhe fshin informacione), atëherë kjo sjellje mund të jetë problematike.
Dhe mati thellësinë e hetimit të tabelës hash pas 64 milionë futjeve (faktori i përdorimit 0,5). Thellësia mesatare ishte 0,4774, kështu që shumica e çelësave ishin ose në slotin më të mirë të mundshëm ose një slot larg pozicionit më të mirë. Thellësia maksimale e tingullit ishte 60.
Më pas mata thellësinë e sondimit në një tryezë me 124 milionë futje (faktori i shfrytëzimit 0,97). Thellësia mesatare ishte tashmë 10,1757, dhe maksimumi - 6474 (!!). Performanca e ndjeshmërisë lineare bie ndjeshëm në shkallë të lartë shfrytëzimi.
Është më mirë ta mbani të ulët shkallën e përdorimit të kësaj tabele hash. Por më pas ne rrisim performancën në kurriz të konsumit të memories. Për fat të mirë, në rastin e çelësave dhe vlerave 32-bit, kjo mund të justifikohet. Nëse në shembullin e mësipërm, në një tabelë me 128 milionë elementë, mbajmë faktorin e shfrytëzimit prej 0,25, atëherë mund të vendosim jo më shumë se 32 milionë elementë në të, dhe 96 milionë lojëra elektronike të mbetura do të humbasin - 8 bajt për çdo palë , 768 MB memorie të humbur.
Ju lutemi vini re se ne po flasim për humbjen e kujtesës së kartës video, e cila është një burim më i vlefshëm se kujtesa e sistemit. Megjithëse shumica e kartave grafike moderne të desktopit që mbështesin CUDA kanë të paktën 4 GB memorie (në momentin e shkrimit, NVIDIA 2080 Ti ka 11 GB), megjithatë nuk do të ishte vendimi më i mençur për të humbur sasi të tilla.
Më vonë do të shkruaj më shumë rreth krijimit të tabelave hash për kartat video që nuk kanë probleme me thellësinë e kërkimit, si dhe mënyrat për të ripërdorur lojëra elektronike të fshira.
Matja e thellësisë së tingullit
Për të përcaktuar thellësinë e kërkimit të një çelësi, ne mund të nxjerrim hash-in e çelësit (indeksin e tij ideal të tabelës) nga indeksi aktual i tabelës:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Për shkak të magjisë së numrave binar të plotësimit të dy-ve dhe faktit që kapaciteti i tabelës hash është dy në fuqinë e dy, kjo qasje do të funksionojë edhe kur indeksi kryesor zhvendoset në fillim të tabelës. Le të marrim një çelës që hash në 1, por është futur në folenë 3. Më pas për një tabelë me kapacitet 4 marrim (3 — 1) & 3, e cila është e barabartë me 2.
Përfundim
Nëse keni pyetje ose komente, ju lutemi më dërgoni email në ose hapni një temë të re në.
Ky kod është shkruar nën frymëzim nga artikuj të shkëlqyer:
Në të ardhmen, do të vazhdoj të shkruaj për zbatimin e tabelave hash për kartat video dhe të analizoj performancën e tyre. Planet e mia përfshijnë lidhjen me zinxhirë, hashimin e Robin Hood dhe hashimin e qyqeve duke përdorur operacione atomike në strukturat e të dhënave që janë miqësore me GPU.
Burimi: www.habr.com