E publikova në Github .
Ky është një tabelë hash e thjeshtë për GPU, e aftë të përpunojë qindra milion insertesh në sekondë. Në laptopin tim me NVIDIA GTX 1060, kodi inserton 64 milion palë çelësa-vlera të gjeneruara rastësisht për rreth 210 ms dhe shemb 32 milion palë për rreth 64 ms.
Pra, shpejtësia në laptop është rreth 300 milion insertesh/s dhe 500 milion shembjesh/s.
Tabela është shkruar në CUDA, ndonëse i njëjti metodë mund të aplikohet në HLSL ose GLSL. Implementimi ka disa kufij që sigurojnë performancë të lartë në kartën grafoike:
- Përpunohen vetëm çelësa dhe vlera 32-bit.
- Tabela hash ka një madhësi fikse.
- Dhe kjo madhësi duhet të jetë e barabartë me dy në fuqinë.
Për çelësat dhe vlerat, duhet të rezervoni një ndarës të thjeshtë (në kodin e dhënë, ky është 0xffffffff).
Tabela hash pa bllokime
Tabela hash përdor adresimin e hapur me , pra është thjesht një array çelës-vlerë që ruhet në memorie dhe ka performancë të shkëlqyer në cache. Kjo nuk është e vërtetë për lidhjen me zinxhir (chaining), e cila përfshin kërkimin e një treguesi në një listë të lidhur. Tabela hash është një array i thjeshtë që ruan elementet KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Madhësia e tabelës është një fuqi e dy, jo një numër i thjeshtë, sepse për të aplikuar maskën pow2/AND mjafton një instruktion e shpejtë, ndërsa operatori modulu punon shumë më ngadalë. Kjo është e rëndësishme në rastin e sondimit linear, pasi gjatë kërkimeve lineare në tabelë, indeksi i slot-it duhet të përkthehet në çdo slot. Dhe si rezultat, shtohet kostoja e operacionit të modulos në çdo slot.
Tabela ruan vetëm çelësin dhe vlerën për secilin element, jo hash-in e çelësit. Duke qenë se tabela ruan vetëm çelësa 32-bit, hash-i llogaritet shumë shpejt. Në kodin e dhënë përdoret hash-i Murmur3, i cili kryen vetëm disa kthesa, XOR dhe shumëzime.
Në tabelë hash aplikohet një metodë mbrojtëse nga bllokimet, e cila nuk varet nga renditja e vendosjes në memorie. Edhe nëse disa operacione shkrimi shkelin renditjen e operacioneve të tjera të tilla, tabela hash ende do të mbajë një gjendje të saktë. Këtë do ta diskutojmë më poshtë. Metoda funksionon shkëlqyeshëm me kartat grafoike, ku kryhen me qindra rrjedha.
ĂelĂ«sat dhe vlerat nĂ« tabelĂ«n hash inicializohen tĂ« zbrazĂ«ta.
Kodi mund të modifikohet për të përpunuar edhe çelësa dhe vlera 64-bit. Për çelësat kërkohen operacione atomike të leximit, shkrimit dhe krahasimit me shkëmbim (compare-and-swap). Ndërsa për vlerat kërkohen operacione atomike të leximit dhe shkrimit. Fatmirësisht, në CUDA operacionet e leximit-shkrimit për vlera 32- dhe 64-bit janë atomike për aq sa ato janë të rreshtuara natyrshëm (shiko ), dhe kartat grafoike moderne mbështesin operacione atomike 64-bit për krahasim me shkëmbim. Sigurisht, gjatë kalimit në 64-bit, performanca do të ulet pak.
Gjendja e tabelës hash
Ădo çift çelĂ«s-vlerĂ« nĂ« tabelĂ«n hash mund tĂ« ketĂ« njĂ« nga katĂ«r gjendje:
- ĂelĂ«si dhe vlera janĂ« tĂ« zbrazĂ«ta. NĂ« kĂ«tĂ« gjendje, tabela hash inicializohet.
- ĂelĂ«si Ă«shtĂ« shkruar, por vlera ende jo. NĂ«se njĂ« rrjedhĂ« tjetĂ«r e ekzekutimit lexon tĂ« dhĂ«nat nĂ« kĂ«tĂ« moment, atĂ«herĂ« ajo kthen njĂ« vlerĂ« tĂ« zbrazĂ«t. Kjo Ă«shtĂ« e zakonshme, do tĂ« ndodhte e njĂ«jta gjĂ« nĂ«se njĂ« rrjedhĂ« tjetĂ«r e ekzekutimit do tĂ« kishte pĂ«rfunduar pak mĂ« herĂ«t, dhe ne flasim pĂ«r njĂ« strukturĂ« tĂ« dhĂ«nash konkurruese.
- Të dy çelësi dhe vlera janë shkruar.
- Vlera është e disponueshme për rrjedhat e tjera të ekzekutimit, ndërsa çelësi ende jo. Kjo mund të ndodhte sepse modeli i programimit në CUDA nënkupton një model të rendit të dobët të memories. Kjo është normale, në çdo rast çelësi ende mbetet i zbrazët, edhe nëse vlera nuk është më e tillë.
NjĂ« detaj i rĂ«ndĂ«sishĂ«m Ă«shtĂ« se sapo çelĂ«si Ă«shtĂ« shkruar nĂ« slot, ai nuk zhvendoset mĂ« â madje edhe nĂ«se çelĂ«si fshihet, pĂ«r kĂ«tĂ« do tĂ« flasim mĂ« poshtĂ«.
Kodi i tabelës hash funksionon edhe me modele të memories me rendit të dobët, në të cilat nuk dihet rendi i leximeve dhe shkrimeve në memorie. Kur të shqyrtojmë inserimin, kërkimin dhe fshirjen në tabelën hash, mbani mend se çdo çift çelës-vlerë ndodhet në njërën nga katër gjendjet e përshkruara më sipër.
Inserimi në tabelën hash
Funksioni CUDA që inserton çifte çelës-vlerë në tabelën hash duket kështu:
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 iteron mbi array-n e tabelës së heshteve duke filluar nga hesi i çelësit të futur. Në çdo vend të array-t, kryhet një operacion atomar krahasimi me ndërrim, ku çelësi në atë vend krahasohet me të zbrazët. Nëse zbulohet një mos përputhje, çelësi në vendin përditësohet me çelësin e futur, dhe më pas kthehet çelësi origjinal i vendit. Nëse ky çelës origjinal ishte i zbrazët ose përputhej me çelësin e futur, atëherë kodi ka gjetur një vend të përshtatshëm për t'u futur dhe vendos vlerën e futur aty.
Në një thirrje të vetme të bërthamës gpu_hashtable_insert() Nëse ka disa elemente me të njëjtin çelës, atëherë çdo nga vlerat e tyre mund të shkruhet në vendin e çelësit. Kjo konsiderohet normale: një nga operacionet e shkruar të çelësit-vlerë gjatë thirrjes do të jetë e suksesshme, por pasi gjithçka ndodh paralelisht brenda disa thread-eve të ekzekutimit, ne nuk mund ta parashikojmë se cili operacion i shkruan në memorie do të jetë i fundit.
Kërkimi në tabelën e heshteve
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 që ruhet në tabelë, ne iteron mbi array-n duke filluar nga hesi i çelësit të kërkuar. Në çdo vend, kontrollojmë nëse çelësi është ai që kërkojmë, dhe nëse po, kthejmë vlerën e tij. Po ashtu, kontrollojmë nëse çelësi është i zbrazët, dhe nëse po, ndalojmë kërkimin.
Nëse nuk arrijmë të gjejmë çelësin, atëherë kodi kthen vlerën e zbrazët.
TĂ« gjitha kĂ«to operacione kĂ«rkimi mund tĂ« kryhen nĂ« mĂ«nyrĂ« pĂ«rmirĂ«suese gjatĂ« futjeve dhe fshirjeve. Ădo çift nĂ« tabelĂ« do tĂ« ketĂ« pĂ«r thread-in njĂ« nga katĂ«r gjendjet e pĂ«rshkruara mĂ« lart.
Fshirja në tabelën e heshteve
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 çelësit realizohet në mënyrë të pazakontë: ne e lëmë çelësin në tabelë dhe e shënojmë vlerën e tij (jo vetë çelësin) si të zbrazët. Ky kod është shumë i ngjashëm me lookup(), përveç se kur zbulohet një përputhje për çelësin, e bën vlerën e tij të zbrazët.
Siç u pĂ«rmend mĂ« sipĂ«r, sapo çelĂ«si shkruhet nĂ« vend, ai nuk lĂ«viz mĂ«. Edhe kur fshihet njĂ« element nga tabela, çelĂ«si mbetet nĂ« vend, thjesht vlera e tij bĂ«het e zbrazĂ«t. Kjo do tĂ« thotĂ« se nuk kemi nevojĂ« pĂ«r tĂ« pĂ«rdorur njĂ« operacion atomar pĂ«r shkrimin e vlerĂ«s sĂ« vendit, sepse nuk ka rĂ«ndĂ«si nĂ«se vlera aktuale Ă«shtĂ« e zbrazĂ«t apo jo â ajo gjithsesi do tĂ« bĂ«het e zbrazĂ«t.
Ndryshimi i madhësisë së tabelës së heshteve
Mund të ndryshoni madhësinë e tabelës së heshteve duke krijuar një tabelë më të madhe dhe duke futur elementët e paplotësuar nga tabela e vjetër në të. Unë këtë funksionalitet nuk e kam realizuar, sepse desha ta mbaja kodin të thjeshtë. Më tej, në programet CUDA, alokimi i memories shpesh bëhet në kodin host, dhe jo në bërthamën CUDA.
Në artikullin përshkruan si të ndryshoni një strukturë të tillë të dhënash, që është e mbrojtur nga bllokimi.
Konkurrenca
Në fragmentet më sipër, funksionet gpu_hashtable_insert(), _lookup() dhe _delete() përpunojnë një çift çelës-vlerë në një kohë. Ndërsa më poshtë gpu_hashtable_insert(), _lookup() dhe _delete() përpunojnë një array çiftesh paralelisht, çdo çift në një thread të veçantë të GPU-së:
// 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 e heshteve me mbrojtje nga bllokimi mbështet futje, kërkime dhe fshirje konkurrente. Duke qenë se çiftet çelës-vlerë gjithmonë ndodhen në njërën nga katër gjendjet, dhe çelësat nuk lëvizin, tabela garanton saktësinë edhe kur përdoren operacione të llojeve të ndryshme njëherësh.
MegjithatĂ«, nĂ«se ne pĂ«rpunojmĂ« paralelisht njĂ« paketĂ« futjesh dhe fshirjesh, dhe nĂ«se nĂ« array-n hyrĂ«s ka çelĂ«sa tĂ« pĂ«rsĂ«ritur, atĂ«herĂ« ne nuk mund tĂ« parashikojmĂ« se cilat çifte "do tĂ« fitojnĂ«" â do tĂ« shkruhen si çifte tĂ« fundit nĂ« tabelĂ«n e heshteve. Supozoni se kemi thirrur kodin e futjes me njĂ« array hyrĂ«s çiftesh A/0 B/1 A/2 C/3 A/4. Kur kodi tĂ« pĂ«rfundojĂ«, çiftet B/1 dhe C/3 do tĂ« jenĂ« sigurisht nĂ« tabelĂ«, ndĂ«rsa aty do tĂ« ndodhet çdo nga çiftet A/0, A/2 ose A/4. Kjo mund tĂ« jetĂ« njĂ« problem, ose ndoshta jo â gjithçka varet nga pĂ«rdorimi. Ju mund tĂ« dini paraprakisht se nĂ« array-n hyrĂ«s nuk ka çelĂ«sa tĂ« pĂ«rsĂ«ritur, ose mund tĂ« mos jetĂ« e rĂ«ndĂ«sishme pĂ«r ju se cila vlerĂ« u regjistrua e fundit.
Nëse kjo është një problem për ju, duhet të ndahen çiftet e dyfishta në thirrje të ndryshme CUDA. Në CUDA, çdo operacion me thirrje bërthamash përfundon përpara thirrjes tjetër të bërthamës (të paktën brenda një rrjedhe. Në rrjedha të ndryshme bërthamat ekzekutohen paralelisht). Në shembullin e mësipërm, nëse thërrisni një bërthamë me A/0 B/1 A/2 C/3, ndërsa një tjetër me A/4, atëherë çelësi A do të marrë vlerën 4.
Tani le të flasim nëse funksionet lookup() dhe delete() duhet të përdorin një tregues të thjeshtë (plain) apo një tregues të ndryshueshëm (volatile) në një masë çiftësh në tabelën e hash-it. firmos se:
Kompilatori mund të optimizojë operacionet e leximit dhe shkruan në memorien globale ose të përbashkët sipas dëshirsë... Këto optimizime mund të çaktivizohen me fjalën kyçe
volatile:... çdo referim në këtë variabël kompilohen në një instruksion të vërtetë leximi ose shkruani në memorie.
Konsideratat e saktësisë nuk kërkojnë aplikimin e saj. volatileNëse rrjedha e ekzekutimit përdor një vlerë të ruajtur nga një operacion më herët të leximit, kjo do të thotë se do të përdorë informacion të pakësuar disi. Por akoma, ky është informacion nga një gjendje e saktë e tabelës së hash-it në një moment të caktuar të thirrjes së bërthamës. Nëse keni nevojë të përdorni informacionin më të freskët, atëherë mund të përdorni një tregues volatile, por atëherë pak do të ulet performanca: sipas testeve të mia, kur hiqen 32 milion elementë, shpejtësia ra nga 500 milion fshirjesh/sec në 450 milion fshirjesh/sec.
Performanca
Në testin e futjes së 64 milion elementëve dhe fshirjes së 32 milion prej tyre, konkurenca mes std::unordered_map dhe tabelës së hash-it për GPU, në fakt nuk ka qenë:
std::unordered_map ka shpenzuar 70 691 ms për të futur dhe fshirë elementët me një çlirim të mëpasshëm unordered_map (çlirimi nga miliona elementë kërkon një kohë të konsiderueshme, sepse brenda unordered_map kryhen shumë alokime memorjeje). Sinqerisht, std:unordered_map ka kufizime shumë të tjera. Kjo është një rrjedhë ekzekutimi e vetme CPU, ajo mbështet çelësa-vlera të çdo madhësie, funksionon mirë me koeficientë të lartë përdorimi dhe tregon performancë të stabilizuar pas shumë fshirjeve.
Koha e funksionimit të tabelës së hash-it për GPU dhe ndërveprimin ndërprogramor ishte 984 ms. Këtu përfshihet koha e shpenzuar për vendosjen e tabelës në memorie dhe heqjen e saj (një herë alokimi i 1 GB memorjeje, që në CUDA merr një kohë të caktuar), si dhe futja dhe fshirja e elementëve, plus iterimi mbi ta. Të gjitha kopjet në memorien dhe nga memoria e kartës grafike gjithashtu janë marrë parasysh.
Operimi vetë i tabelës së hash-it zuri 271 ms. Këtu përfshihet koha e shpenzuar nga karta grafike për të futur dhe fshirë elementët dhe nuk përfshin kohën për kopjimin në memorie dhe iterimin mbi tabelën e rezultuar. Nëse tabela GPU jeton gjatë, ose nëse tabela e hash-it është plotësisht brenda memorjes së kartës grafike (për shembull, për krijimin e një tabele hash që do të përdoret nga kodi tjetër të GPU dhe jo nga procesori qendror), atëherë rezultati i testimit është relevant.
Tabela e hash-it për kartën grafike tregon performancë të lartë falë gjerësisë së madhe të bandës dhe paralelizimit aktiv.
Disavantazhet
Arkitektura e tabelës së hash-it ka disa probleme që duhet të mbahen parasysh:
- Zbulimi linear pengohet nga grumbullimi, për shkak të cilit çelësat në tabelë nuk vendosen aq mirë.
- ĂelĂ«sat nuk fshihen me funksionin
fshidhe me kalimin e kohës i mbingarkojnë tabelën.
Si rezultat, performanca e tabelës së hash-it mund të zvogëlohet gradualisht, sidomos nëse ajo ekziston për një kohë të gjatë dhe ka shumë futje dhe fshirje. Një nga mënyrat për të lehtësuar këto disavantazhe është reheshimi në një tabelë të re me një koeficient përdorimi mjaft të ulët dhe filtrimi i çelësave të fshirë gjatë reheshimit.
Për të ilustruar problemet e përshkruara, do të përdor kodin e mësipërm për të krijuar një tabelë me 128 milion elementë, duke futur ciklikisht 4 milion elementë, derisa të mbush 124 milion vende (koeficienti i përdorimit rreth 0,96). Këtu është tabela e rezultateve, çdo rresht është një thirrje bërthamë CUDA me futjen e 4 milion elementëve të rinj në një tabelë hash:
Koeficienti i përdorimit
Koha e futjes së 4 194 304 elementëve
0,00
11,608448 ms (361,314798 milion çelësa/sec.)
0,03
11,751424 ms (356,918799 milion çelësa/sec.)
0,06
11,942592 ms (351,205515 milion çelësa/sec.)
0,09
12,081120 ms (347,178429 milion çelësa/sec.)
0,12
12,242560 ms (342,600233 milion çelësa/sec.)
0,16
12,396448 ms (338,347235 milion çelësa/sec.)
0,19
12,533024 ms (334,660176 milion çelësa/sec.)
0,22
12,703328 ms (330,173626 milion çelësa/sec.)
0,25
12,884512 ms (325,530693 milion çelësa/sec.)
0,28
13,033472 ms (321,810182 milion çelësa/sec.)
0,31
13,239296 ms (316,807174 milion çelësa/sec.)
0,34
13,392448 ms (313,184256 milion çelësa/sec.)
0,37
13,624000 ms (307,861434 milion çelësa/sec.)
0,41
13,875520 ms (302,280855 milion çelësa/sec.)
0,44
14,126528 ms (296,909756 milion çelësash/sek.)
0,47
14,399328 ms (291,284699 milion çelësash/sek.)
0,50
14,690304 ms (285,515123 milion çelësash/sek.)
0,53
15,039136 ms (278,892623 milion çelësash/sek.)
0,56
15,478656 ms (270,973402 milion çelësash/sek.)
0,59
15,985664 ms (262,379092 milion çelësash/sek.)
0,62
16,668673 ms (251,627968 milion çelësash/sek.)
0,66
17,587200 ms (238,486174 milion çelësash/sek.)
0,69
18,690048 ms (224,413765 milion çelësash/sek.)
0,72
20,278816 ms (206,831789 milion çelësash/sek.)
0,75
22,545408 ms (186,038058 milion çelësash/sek.)
0,78
26,053312 ms (160,989275 milion çelësash/sek.)
0,81
31,895008 ms (131,503463 milion çelësash/sek.)
0,84
42,103294 ms (99,619378 milion çelësash/sek.)
0,87
61,849056 ms (67,815164 milion çelësash/sek.)
0,90
105,695999 ms (39,682713 milion çelësash/sek.)
0,94
240,204636 ms (17,461378 milion çelësash/sek.)
Me rritjen e koeficientit të përdorimit, performanca ulet. Kjo është e papërshtatshme në shumicën e rasteve. Nëse aplikacioni fut elemente në tabelë dhe më pas i hedh poshtë (p.sh., gjatë numërimit të fjalëve në një libër), atëherë kjo nuk është problem. Por nëse aplikacioni përdor një tabelë hëzhuese me jetëgjatësi të gjate (p.sh., në një redaktues grafik për të ruajtur pjesët e papra të imazheve, kur përdoruesi shpesh fut dhe fshin informacion), atëherë ky qëndrim mund të shkaktojë shqetësim.
Dhe mat të thellësinë e sondimit të tabelës hëzhuese pas 64 milion futjeve (koeficienti i përdorimit 0,5). Thellësia mesatare ishte 0,4774, pra shumica e çelësave ishin ose në njërin nga slotet më të mira të mundshme, ose në një slot nga pozita më e mirë. Thellësia maksimale e sondimit ishte 60.
Më pas matës thellësinë e sondimit në tabelën me 124 milion futje (koeficienti i përdorimit 0,97). Thellësia mesatare ishte tashmë 10,1757, ndërsa maksimale - 6474 (!!). Performanca e sondimit linear bie ndjeshëm me koeficient të lartë përdorimi.
Më mirë është të mbash këtë tabelë hëzhuese me një koeficient të ulët përdorimi. Por atëherë, ne rrisim performancën me koston e konsumit të memories. Fatmirësisht, në rastin e çelësave dhe vlerave 32-bitësh, kjo mund të justifikohet. Në shembullin e mësipërm, në një tabelë prej 128 milion elementesh, duke ruajtur një koeficient përdorimi 0,25, ne mund të vendosim maksimumi 32 milion elemente, ndërsa 96 milion slotet e tjera do të shuhen - 8 byte për çdo çift, 768 MB memorje e humbur.
Kujdes, kjo është humbje memorjeje të kartës grafike, e cila është një burim më i çmuar se memoria sistemike. Edhe pse shumica e kartave moderne të grafikës që mbështesin CUDA kanë të paktën 4 GB memorie (në momentin e shkruar, NVIDIA 2080 Ti ka 11 GB), humbja e këtyre sasi do të jetë shume e mençur.
Më vonë do të shkruaj në detaje mbi krijimin e tabelave hëzhuese për kartat grafike që s'kanë probleme me thellësinë e sondimit, si dhe mbi mënyrat e ripërdorimit të slotëve të fshira.
Matja e thellësisë së sondimit
Për të përcaktuar thellësinë e sondimit të çelësit, ne mund të nxjerrim hashin e çelësit (indeksi ideal në tabelë) nga indeksi i tij real tabelar:
// 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ë dy numrave binarë në kodin e plotë dhe faktit që kapaciteti i tabelës hëzhuese është dy në fuqinë, ky qasje do të funksionojë edhe kur indeksi i çelësit zhvendoset në fillim të tabelës. Le të marrim një çelës që hashizohet në 1, por futet në slot 3. Atëherë për një tabelë me kapacitet 4 do të kemi (3 - 1) & 3, që është ekuivalente me 2.
PĂ«rfundimi
Nëse keni pyetje ose komente, më shkruani në ose hapni një temë të re në.
Ky kod është shkruar nën frymëzimin nga artikuj të shkëlqyer:
Në të ardhmen do të vazhdoj të shkruaj rreth realizimeve të tabelave hëzhuese për kartat grafike dhe do të analizoj performancën e tyre. Në planet e mia janë lidhja në zinxhir, hashizimi i Robin Hood dhe hashizimi i kukullës duke përdorur operacione atomike në struktura të dhënash që janë të përshtatshme për kartat grafike.
Burimi: habr.com
