Min ew li ser Github şand .
Ew tabloyek hash a GPU-ya hêsan e ku karibe di çirkeyê de bi sed mîlyonan insertan bixebite. Li ser laptopa min a NVIDIA GTX 1060, kod di nav 64 ms de 210 mîlyon cotên nirx-nirxa ku bi rengek bêserûber hatine hilberandin têxe nav 32 ms.
Ango leza li ser laptopek bi qasî 300 mîlyon têxe / saniye û 500 mîlyon jêbirin / saniye ye.
Tablo di CUDA de hatî nivîsandin, her çend heman teknîk dikare li HLSL an GLSL jî were sepandin. Pêkanînê çend sînor hene da ku performansa bilind li ser qerta vîdyoyê peyda bike:
- Tenê bişkojkên 32-bit û heman nirx têne pêvajo kirin.
- Tabloya hash xwedî pîvanek sabît e.
- Û ev mezinahî divê bi du hêzan re wekhev be.
Ji bo kilît û nirxan, hûn hewce ne ku nîşankerek veqetandî ya hêsan veqetînin (di koda jorîn de ev 0xffffffff e).
Tabloya Haş bê kilît
Tabloya hash navnîşana vekirî bi kar tîne , ango, ew bi tenê komek ji cotên key-nirxê ye ku di bîranînê de tê hilanîn û performansa cache-ya bilindtir heye. Heman tişt ji bo zincîrkirinê nayê gotin, ku tê de lêgerîna nîşankerek di navnîşek girêdayî de heye. Tabloyek hash komek sade ye ku hêmanan hilîne KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Mezinahiya tabloyê hêzek du ye, ne hejmareke sereke ye, ji ber ku yek rêwerzek bilez bes e ku meriv maskeya pow2/AND bikar bîne, lê operatorê modulusê pir hêdîtir e. Ev di mijara lêgerîna xêzikî de girîng e, ji ber ku di lêgerîna tabloya xêz de pêdivî ye ku navnîşa hêlînê di her hêlînê de were pêçandin. Û di encamê de, lêçûna operasyonê di her hêlînê de modul tê zêdekirin.
Tablo tenê ji bo her hêmanek mifteyê û nirxê hildide, ne haşek mifteyê. Ji ber ku tablo tenê bişkojkên 32-bit hilîne, hash pir zû tê hesibandin. Koda li jor hash Murmur3 bikar tîne, ku tenê çend guheztin, XOR û pirjimaran pêk tîne.
Tabloya hash teknîkên parastinê yên girtinê yên ku ji rêza bîranînê serbixwe ne bikar tîne. Her çend hin operasyonên nivîsandinê rêza operasyonên din ên bi vî rengî xera bikin jî, tabloya hash dê dîsa jî rewşa rast biparêze. Em ê li jêr li ser vê biaxivin. Teknîkî bi kartên vîdyoyê yên ku bi hezaran mijaran bi hev re dimeşînin re pir baş dixebite.
Bişkok û nirxên di tabloya hash de ji bo valakirinê têne destpêkirin.
Kod dikare were guheztin da ku bişkok û nirxên 64-bit jî bigire. Bişkojkên atomî xwendin, nivîsandin, û berawirdkirin-û-swap operasyonên hewce dike. Û nirx xebatên xwendin û nivîsandina atomî hewce dike. Xweşbextane, di CUDA de, operasyonên xwendin-nivîsandinê ji bo nirxên 32- û 64-bit atomî ne heya ku ew bi xwezayî li hevûdu ne (li jêr binêre). ), û kartên vîdyoyê yên nûjen piştgirî didin operasyonên berhevkirin û veguheztina atomî yên 64-bit. Bê guman, dema ku diçin 64 bit, performans dê hinekî kêm bibe.
Rewşa sifrê Hash
Her cotek nirx-kilît di tabloyek hash de dikare yek ji çar rewşan hebe:
- Key û nirx vala ne. Di vê rewşê de, tabloya hash tê destpêkirin.
- Kilîtê hatiye nivîsandin, lê hêj nirx nehatiye nivîsandin. Ger mijarek din niha daneyê dixwîne, wê hingê ew vala vedigere. Ev normal e, heman tişt dê biqewimiya heke mijarek din a darvekirinê hinekî berê bixebitiya, û em li ser avahiyek daneya hevdemî diaxivin.
- Hem key û hem jî nirx têne tomar kirin.
- Nirx ji mijarên din ên darvekirinê re peyda dibe, lê kilît hîn nebûye. Ev dikare biqewime ji ber ku modela bernamesaziya CUDA xwedan modelek bîranînê ya bi rêzkirî ye. Ev normal e; di her rewşê de, mift hîn jî vala ye, her çend nirx êdî ne wusa be.
Nîşanek girîng ev e ku gava ku mift li hêlînê hat nivîsandin, ew nema dimeşe - heke mift were jêbirin jî, em ê li jêr li ser vê biaxivin.
Koda tabloya hash tewra bi modelên bîranînê yên bi rêzkirî re dixebite ku tê de rêza ku bîranîn tê xwendin û nivîsandin nayê zanîn. Gava ku em di tabloyek hash de li danîn, lêgerîn û jêbirinê dinêrin, ji bîr mekin ku her cotek key-nirx di yek ji çar dewletên ku li jor hatine destnîşan kirin de ye.
Têkeve nav tabloyek hash
Fonksiyona CUDA ya ku cotên key-nirxê têxe tabloyek hash wiha xuya dike:
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);
}
}
Ji bo ku mifteyek têxin, kod di nav rêza tabloya hash-ê de ku bi haşeya mifteya têxê dest pê dike tê dubare kirin. Her slotek di rêzê de operasyonek berhevkirin-û-guhartinê ya atomî pêk tîne ku mifteya di wê hêlînê de vala dike. Ger neliheviyek were dîtin, mifteya di hêlînê de bi mifteya têxistî tê nûve kirin, û dûv re mifteya hêlînê ya orîjînal tê vegerandin. Ger ev mifteya orîjînal vala bûya an jî bi mifteya lêxistî re lihevhatî bûya, wê demê kod ji bo têxistinê hêlînek minasib peyda kir û nirxa lêxistî têxiste hêlînê.
Ger di yek kernelê de bang bikin gpu_hashtable_insert() bi heman mifteyê gelek hêman hene, wê hingê yek ji nirxên wan dikare li hêlîna mifteyê were nivîsandin. Ev normal tê hesibandin: yek ji wan key-nirxa ku di dema bangê de dinivîse dê biserkeve, lê ji ber ku ev hemî di nav çend mijarên darvekirinê de paralel diqewime, em nekarin pêşbînî bikin ka kîjan nivîsandina bîranînê dê ya paşîn be.
Lêgerîna tabloya Hash
Koda ji bo mifteyên lêgerînê:
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);
}
}
Ji bo ku em nirxa mifteyek ku di tabloyek de hatî hilanîn bibînin, em di nav rêza ku bi hashê mifteya ku em lê digerin dest pê dike dubare dikin. Di her hêlînê de, em kontrol dikin ka kilît ew e ku em lê digerin, û heke wusa be, em nirxa wê vedigerînin. Em jî kontrol dikin ka mift vala ye, û heke wusa be, em lêgerînê betal dikin.
Ger em nikaribin mifteyê bibînin, kod nirxek vala vedigerîne.
Hemî van operasiyonên lêgerînê dikarin bi navgîn û jêbirinê bi hev re bêne kirin. Her cotek di tabloyê de dê yek ji çar dewletên ku li jor ji bo herikînê hatine destnîşan kirin hebe.
Di tabloyek hash de jêbirin
Koda ji bo jêbirina mifteyan:
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);
}
}
Jêbirina mifteyê bi rengekî neasayî pêk tê: em mifteyê di tabloyê de dihêlin û nirxa wê (ne kilîta bi xwe) wekî vala nîşan didin. Ev kod pir dişibihe lookup(), ji xeynî ku gava lihevkirinek li ser mifteyê tê dîtin, nirxa wê vala dike.
Wekî ku li jor jî hate behs kirin, gava ku mifteyek li hêlekê were nivîsandin, êdî nayê veguheztin. Dema ku hêmanek ji tabloyê were jêbirin jî, mift li cihê xwe dimîne, nirxa wê bi tenê vala dibe. Ev tê wê wateyê ku em ne hewce ne ku ji bo nirxa hêlînê operasyonek nivîsandina atomî bikar bînin, ji ber ku ne girîng e nirxa heyî vala ye an na - ew ê dîsa jî vala bibe.
Mezinahiya tabloyek hash
Hûn dikarin bi çêkirina tabloyek mezintir û xistina hêmanên ne vala ji tabloya kevn têxin mezinahiya tabloyek hash biguherînin. Min vê fonksiyonê bicîh neanî ji ber ku min dixwest ku koda nimûne sade bihêlim. Digel vê yekê, di bernameyên CUDA de, veqetandina bîranînê bi gelemperî di koda mêvandar de ji kernel CUDA tête kirin.
Di gotarê de diyar dike ka meriv çawa avahiyek daneya parastî ya bi vî rengî biguhezîne.
Competitiveness
Di perçeyên koda fonksiyona jorîn de gpu_hashtable_insert(), _lookup() и _delete() di yek carê de cotek key-nirxê pêvajoyê bikin. Û kêmtir gpu_hashtable_insert(), _lookup() и _delete() Rêzeya cotan bi paralelî pêvajoyê bikin, her cotek di mijarek cîbecîkirina GPU-yê de:
// 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);
}
}
Tabloya hash-berxwedêr a kilîtkirî piştgirî, lêgerandin û jêbirinên hevdemî piştgirî dike. Ji ber ku cotên key-nirx her gav di yek ji çar rewşan de ne û bişkok nagerin, tablo rastbûna xwe garantî dike hetta dema ku celebên cûda operasyonên hevdem têne bikar anîn.
Lêbelê, heke em komek têxistin û jêbirinan bi paralelî pêvajoyê bikin, û heke di rêza têketinê ya cotan de mifteyên dubare hebin, wê hingê em ê nikaribin pêşbînî bikin ka kîjan cotên dê "biserkevin" - dê herî dawî li ser tabloya haş were nivîsandin. Ka em bibêjin ku me koda têxê bi komek ketina cotan re gazî kir A/0 B/1 A/2 C/3 A/4. Dema ku kod biqede, cotên B/1 и C/3 garantî kirin ku di tabloyê de hebin, lê di heman demê de yek ji wan cotan dê tê de xuya bibe A/0, A/2 an A/4. Dibe ku ev pirsgirêk be an jî nebe - ew hemî bi serîlêdanê ve girêdayî ye. Dibe ku hûn ji berê de zanibin ku di rêzika têketinê de bişkokên dubare nînin, an jî dibe ku hûn bala xwe nedin kîjan nirxê paşîn hatî nivîsandin.
Ger ev ji bo we pirsgirêkek e, wê hingê hûn hewce ne ku cotên dubare di nav bangên pergala CUDA yên cihêreng de veqetînin. Di CUDA de, her operasiyona ku gazî kernelê dike, her gav berî banga kernelê ya din temam dibe (kêmtir di nav mijarekê de. Di mijarên cûda de, kernel bi hev re têne darve kirin). Di mînaka li jor de, heke hûn bi yek kernelê re dibêjin A/0 B/1 A/2 C/3, û ya din bi A/4, paşê key A dê qîmetê bigire 4.
Naha em li ser bipeyivin ka fonksiyon divê an na lookup() и delete() Di tabloya haş de nîşanek sade an guhezbar bikar bînin. Dewletên ku:
Berhevkar dikare bibijêre ku xwendin û nivîsandinê li bîra gerdûnî an hevbeş xweşbîn bike... Van xweşbînkirin dikarin bi karanîna peyva sereke werin neçalak kirin
volatile: ... her referansek ji bo vê guhêrbar di nav fermanek xwendin an nivîsandinê ya bîranîna rastîn de tê berhev kirin.
Fikrên rastdariyê serîlêdanê hewce nake volatile. Ger mijara înfazê nirxek cached ji xebatek xwendina berê bikar bîne, wê hingê ew ê agahdariya piçûktir bikar bîne. Lê dîsa jî, ev agahdarî ji rewşa rast a tabloya hash di demek diyarkirî ya banga kernel de ye. Heke hûn hewce ne ku agahdariya herî dawî bikar bînin, hûn dikarin navnîşê bikar bînin volatile, lê wê hingê performans dê hinekî kêm bibe: li gorî ceribandinên min, dema ku 32 mîlyon hêman jêbirin, leza ji 500 mîlyon jêbirin / saniyê daket 450 mîlyon jêbirin / saniyê.
Berhemdariyê
Di îmtîhanê de ji bo têxistina 64 mîlyon hêman û jêbirina 32 mîlyon ji wan, pêşbaziya di navbera std::unordered_map û ji bo GPU bi rastî tabloyek hash tune:
std::unordered_map 70 ms derbas kirin û hêmanan derxistin û paşê wan azad kirin unordered_map (Rabûna ji mîlyon hêmanan gelek dem digire, ji ber ku di hundurê de unordered_map gelek veqetandinên bîranînê têne çêkirin). Bi dilsozî, std:unordered_map sînorkirinên bi temamî cuda. Ew yek Mijara CPU-yê ya darvekirinê ye, nirx-kilîdên bi her mezinahî piştgirî dike, di rêjeyên karanîna bilind de baş çêdibe, û piştî gelek jêbirinan performansa aram nîşan dide.
Demjimêra tabloya hash ji bo GPU û pêwendiya navbernameyê 984 ms bû. Di nav vê de dema danîna tabloyê di bîrê de û jêbirina wê (veqetandina 1 GB bîra yek carê, ku di CUDA de hin dem digire), têxistin û jêbirina hêmanan, û dubarekirina li ser wan vedihewîne. Hemî kopiyên ji bîra qerta vîdyoyê û ji bîrê jî têne hesibandin.
Tabloya hash bixwe 271 ms girt ku temam bibe. Ev dema ku ji hêla qerta vîdyoyê ve tê veqetandin û jêbirina hêmanan vedihewîne, û dema kopîkirina nav bîranînê û dubarekirina li ser tabloya encam de derbas dibe, hesab nake. Ger tabloya GPU demek dirêj bijî, an heke tabloya hash bi tevahî di bîranîna qerta vîdyoyê de hebe (mînakî, ji bo afirandina tabloyek hash ku dê ji hêla koda GPU-ya din ve û ne pêvajoya navendî ve were bikar anîn), wê hingê encama testê têkildar e.
Tabloya hash a ji bo qerta vîdyoyê ji ber berbi bilind û paralelbûna çalak performansa bilind nîşan dide.
kêmasiyên
Mîmariya tabloya hash çend pirsgirêk hene ku divê ji wan haydar bin:
- Lêpirsîna xêzkirî ji hêla kombûnê ve tê asteng kirin, ku dibe sedem ku mifteyên di tabloyê de ji bêkêmasî kêmtir werin danîn.
- Keys bi karanîna fonksiyonê nayê rakirin
deleteû bi demê re ew sifrê tevlihev dikin.
Wekî encamek, performansa tabloyek hash dikare hêdî hêdî kêm bibe, nemaze heke ew ji bo demek dirêj ve hebe û gelek têkel û jêbirin hebin. Yek rê ji bo sivikkirina van kêmasiyan ev e ku meriv di tabloyek nû de bi rêjeyek karanîna pir kêm veguhezîne û bişkokên jêkirî di dema rehashkirinê de fîlter bike.
Ji bo ronîkirina pirsgirêkên ku hatine destnîşan kirin, ez ê koda jorîn bikar bînim da ku tabloyek bi 128 mîlyon hêmanan biafirînim û di nav 4 mîlyon hêmanan de bigerim heya ku ez 124 mîlyon hêlîn dagirim (rêjeya karanîna bi qasî 0,96). Li vir tabloya encamê ye, her rêz bangek kernelê CUDA ye ku 4 mîlyon hêmanên nû têxe tabloyek haş:
Rêjeya Bikaranîna
Demjimêra têketina 4 hêmanan
0,00
11,608448 ms (361,314798 mîlyon kilît / saniye.)
0,03
11,751424 ms (356,918799 mîlyon kilît / saniye.)
0,06
11,942592 ms (351,205515 mîlyon kilît / saniye.)
0,09
12,081120 ms (347,178429 mîlyon kilît / saniye.)
0,12
12,242560 ms (342,600233 mîlyon kilît / saniye.)
0,16
12,396448 ms (338,347235 mîlyon kilît / saniye.)
0,19
12,533024 ms (334,660176 mîlyon kilît / saniye.)
0,22
12,703328 ms (330,173626 mîlyon kilît / saniye.)
0,25
12,884512 ms (325,530693 mîlyon kilît / saniye.)
0,28
13,033472 ms (321,810182 mîlyon kilît / saniye.)
0,31
13,239296 ms (316,807174 mîlyon kilît / saniye.)
0,34
13,392448 ms (313,184256 mîlyon kilît / saniye.)
0,37
13,624000 ms (307,861434 mîlyon kilît / saniye.)
0,41
13,875520 ms (302,280855 mîlyon kilît / saniye.)
0,44
14,126528 ms (296,909756 mîlyon kilît / saniye.)
0,47
14,399328 ms (291,284699 mîlyon kilît / saniye.)
0,50
14,690304 ms (285,515123 mîlyon kilît / saniye.)
0,53
15,039136 ms (278,892623 mîlyon kilît / saniye.)
0,56
15,478656 ms (270,973402 mîlyon kilît / saniye.)
0,59
15,985664 ms (262,379092 mîlyon kilît / saniye.)
0,62
16,668673 ms (251,627968 mîlyon kilît / saniye.)
0,66
17,587200 ms (238,486174 mîlyon kilît / saniye.)
0,69
18,690048 ms (224,413765 mîlyon kilît / saniye.)
0,72
20,278816 ms (206,831789 mîlyon kilît / saniye.)
0,75
22,545408 ms (186,038058 mîlyon kilît / saniye.)
0,78
26,053312 ms (160,989275 mîlyon kilît / saniye.)
0,81
31,895008 ms (131,503463 mîlyon kilît / saniye.)
0,84
42,103294 ms (99,619378 mîlyon kilît / saniye.)
0,87
61,849056 ms (67,815164 mîlyon kilît / saniye.)
0,90
105,695999 ms (39,682713 mîlyon kilît / saniye.)
0,94
240,204636 ms (17,461378 mîlyon kilît / saniye.)
Her ku karanîna zêde dibe, performansa kêm dibe. Ev di pir rewşan de nayê xwestin. Ger serîlêdan hêmanan têxe tabloyekê û dûv re wan ji holê rake (mînak, dema ku peyvan di pirtûkekê de dihejmêre), wê hingê ev ne pirsgirêk e. Lê ger serîlêdan tabloyek hash-a-dirêj bikar bîne (mînak, di edîtorek grafîkê de da ku beşên ne vala yên wêneyan hilîne ku bikarhêner bi gelemperî agahdarî lê dixe û jê dike), wê hingê ev tevger dikare bibe pirsgirêk.
Û kûrahiya lêkolîna tabloya hash piştî 64 mîlyon têketinê pîva (faktora karanîna 0,5). Kûrahiya navîn 0,4774 bû, ji ber vê yekê piraniya kilîtan an di hêlîna çêtirîn gengaz de bûn an jî hêlînek ji pozîsyona çêtirîn dûr bûn. Kûrahiya dengê herî zêde 60 bû.
Dûv re min kûrahiya lêkolînê li ser maseyek bi 124 mîlyon pêvekirî (faktora karanîna 0,97) pîva. Kûrahiya navîn berê 10,1757 bû, û herî zêde - 6474 (!!). Performansa hestiyariya xêzkirî di rêjeyên karanîna bilind de pir kêm dibe.
Ya çêtirîn e ku meriv rêjeya karanîna vê tabloya hash kêm bimîne. Lê dûv re em performansê li ser xerckirina bîranînê zêde dikin. Xweşbextane, di doza bişkok û nirxên 32-bit de, ev dikare rastdar be. Ger di mînaka jorîn de, di tabloyek bi 128 mîlyon hêmanan de, em faktora karanîna 0,25 biparêzin, wê hingê em nikanin ji 32 mîlyon hêmanan bêtir tê de bi cîh bikin, û 96 mîlyon hêlên mayî dê winda bibin - ji bo her cotek 8 byte. , 768 MB bîra wenda.
Ji kerema xwe not bikin ku em behsa windakirina bîranîna qerta vîdyoyê dikin, ku ji bîra pergalê çavkaniyek hêjatir e. Her çend pir kartên grafikên sermaseyê yên nûjen ên ku CUDA piştgirî dikin bi kêmî ve 4 GB bîranîn hene (di dema nivîsandinê de, NVIDIA 2080 Ti xwedan 11 GB ye), dîsa jî ew ê ne biryara herî aqilmend be ku meriv mîqdarên weha winda bike.
Dûv re ez ê bêtir li ser çêkirina tabloyên hash ji bo kartên vîdyoyê yên ku di kûrahiya lêkolînê de pirsgirêk nînin, û her weha awayên ji nû ve karanîna hêlînên jêbirin de binivîsim.
Pîvana kûrahiya deng
Ji bo destnîşankirina kûrahiya lêgerîna mifteyê, em dikarin kêşeya mifteyê (nîşana tabloya wê ya îdeal) ji navnîşa wê ya tabloya rastîn derxînin:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Ji ber sêrbaziya hejmarên binaryê yên temamkerê du-duyan û ji ber ku kapasîteya tabloya hash du bi hêza duyan e, ev nêzîkatî dê bixebite tewra gava ku nîşaneya sereke berbi destpêka tabloyê ve were veguheztin. Werin em kilîtek ku bi 1-ê ve hatî heşe kirin, lê têxe hêlîna 3-yê. Dûv re ji bo tabloyek bi kapasîteya 4 em digirin (3 — 1) & 3, ku bi 2 re wekhev e.
encamê
Ger pirs an şîroveyên we hebin, ji kerema xwe ji min re bi e-nameyê re bişînin an jî mijarek nû vekin.
Ev kod di bin îlhama gotarên hêja de hate nivîsandin:
Di pêşerojê de, ez ê berdewam bikim ku li ser pêkanînên tabloya hash-ê ji bo kartên vîdyoyê binivîsim û performansa wan analîz bikim. Di planên min de zincîrkirin, çewisandina Robin Hood, û xerakirina cuckoo bi karanîna operasyonên atomî di strukturên daneyê yên ku GPU-yê de ne, hene.
Source: www.habr.com