Mwen poste li sou Github .
Li se yon senp GPU hash tab ki kapab trete dè santèn de milyon foure pou chak segonn. Sou laptop NVIDIA GTX 1060 mwen an, kòd la mete 64 milyon pè kle-valè ki te pwodwi owaza nan apeprè 210 ms epi retire 32 milyon pè nan apeprè 64 ms.
Sa vle di, vitès la sou yon laptop se apeprè 300 milyon foure / sec ak 500 milyon dola efase / sec.
Tablo a ekri nan CUDA, byenke menm teknik la ka aplike nan HLSL oswa GLSL. Aplikasyon an gen plizyè limit pou asire pèfòmans segondè sou yon kat videyo:
- Se sèlman kle 32-bit ak menm valè yo trete.
- Tablo hash la gen yon gwosè fiks.
- Ak gwosè sa a dwe egal a de pouvwa a.
Pou kle ak valè, ou bezwen rezève yon makè delimiter senp (nan kòd ki anwo a sa a se 0xffffffff).
Hash tab san kadna
Tablo hash la itilize adrès louvri ak , se sa ki, li se tou senpleman yon etalaj de pè kle-valè ki estoke nan memwa epi ki gen pèfòmans siperyè kachèt. Menm bagay la tou pa ka di pou chaining, ki enplike rechèch pou yon konsèy nan yon lis lye. Yon tab hash se yon senp etalaj ki estoke eleman KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Gwosè tab la se yon pouvwa de, pa yon nimewo premye, paske yon sèl enstriksyon rapid ase pou aplike mask la pow2/AND, men operatè modil la pi dousman. Sa a enpòtan nan ka a nan sonde lineyè, depi nan yon rechèch tab lineyè endèks la plas dwe vlope nan chak plas. Epi kòm yon rezilta, pri a nan operasyon an ajoute modulo nan chak plas.
Tablo a sèlman magazen kle a ak valè pou chak eleman, pa yon hash nan kle a. Depi tab la sèlman magazen kle 32-bit, hash la kalkile trè vit. Kòd ki pi wo a sèvi ak hash Murmur3, ki sèlman fè kèk chanjman, XOR ak miltiplikasyon.
Tablo hash la itilize teknik pwoteksyon bloke ki endepandan de lòd memwa. Menm si kèk operasyon ekri deranje lòd lòt operasyon sa yo, tab hash la ap toujou kenbe eta kòrèk la. Nou pral pale sou sa anba a. Teknik la travay anpil ak kat videyo ki kouri dè milye de fil ansanm.
Kle yo ak valè yo nan tablo hash yo inisyalize yo vid.
Kòd la ka modifye pou okipe kle 64-bit ak valè tou. Kle mande pou operasyon atomik lekti, ekri, ak konpare ak swap. Ak valè yo mande pou operasyon li ak ekri atomik. Erezman, nan CUDA, operasyon lekti-ekri pou valè 32- ak 64-bit yo atomik osi lontan ke yo natirèlman aliyen (gade anba a). ), ak kat videyo modèn sipòte operasyon konpare ak echanj atomik 64-bit. Natirèlman, lè w ap deplase nan 64 Bits, pèfòmans ap diminye yon ti kras.
Eta tab hash
Chak pè kle-valè nan yon tablo hash ka gen youn nan kat eta:
- Kle ak valè yo vid. Nan eta sa a, tab hash la inisyalize.
- Yo te ekri kle a, men valè a poko ekri. Si yon lòt fil ap li done kounye a, li retounen vid. Sa a se nòmal, menm bagay la ta rive si yon lòt fil nan ekzekisyon te travay yon ti kras pi bonè, epi nou ap pale de yon estrikti done konkouran.
- Tou de kle a ak valè a anrejistre.
- Valè a disponib nan lòt fil nan ekzekisyon, men kle a poko. Sa a ka rive paske modèl pwogramasyon CUDA a gen yon modèl memwa ki pa gen anpil lòd. Sa a nòmal; nan nenpòt ka, kle a toujou vid, menm si valè a pa konsa ankò.
Yon nuans enpòtan se ke yon fwa yo te ekri kle a nan plas la, li pa deplase ankò - menm si kle a efase, nou pral pale sou sa a anba a.
Kòd tab hash la menm travay ak modèl memwa ki pa gen anpil lòd nan ki lòd yo li ak ekri memwa se enkoni. Pandan n ap gade ensèsyon, rechèch, ak sipresyon nan yon tablo hash, sonje ke chak pè kle-valè se nan youn nan kat eta ki dekri pi wo a.
Mete nan yon tab hash
Fonksyon CUDA ki mete pè kle-valè nan yon tablo hash sanble sa a:
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);
}
}
Pou mete yon kle, kòd la repete atravè etalaj tab hash la kòmanse ak hash kle ki mete a. Chak plas nan etalaj la fè yon operasyon konpare ak swap atomik ki konpare kle nan plas sa a ak vid. Si yo detekte yon dezakò, kle nan plas la mete ajou ak kle ki antre, epi apre kle plas orijinal la retounen. Si kle orijinal sa a te vid oswa matche ak kle a antre, Lè sa a, kòd la te jwenn yon plas apwopriye pou ensèsyon ak mete valè a antre nan plas la.
Si nan yon sèl apèl nwayo gpu_hashtable_insert() gen plizyè eleman ak menm kle a, Lè sa a, nenpòt nan valè yo ka ekri nan plas kle a. Sa a konsidere kòm nòmal: youn nan ekri kle-valè yo pandan apèl la pral reyisi, men depi tout bagay sa yo rive nan paralèl nan plizyè fil nan ekzekisyon, nou pa ka predi ki ekri memwa ki pral dènye a.
Fè rechèch tab hash
Kòd pou chèche kle:
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);
}
}
Pou jwenn valè a nan yon kle ki estoke nan yon tablo, nou repete nan etalaj la kòmanse ak hash nan kle nou ap chèche a. Nan chak plas, nou tcheke si kle a se youn nan nou ap chèche a, epi si se konsa, nou retounen valè li yo. Nou tcheke tou si kle a vid, epi si se konsa, nou avòte rechèch la.
Si nou pa ka jwenn kle a, kòd la retounen yon valè vid.
Tout operasyon rechèch sa yo ka fèt ansanm atravè ensèsyon ak sipresyon. Chak pè nan tablo a pral gen youn nan kat eta ki dekri pi wo a pou koule a.
Efase nan yon tab hash
Kòd pou efase kle yo:
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);
}
}
Efase yon kle se yon fason ki pa nòmal: nou kite kle a nan tablo a epi make valè li (pa kle nan tèt li) kòm vid. Kòd sa a sanble anpil ak lookup(), eksepte ke lè yo jwenn yon match sou yon kle, li fè valè li vid.
Kòm mansyone pi wo a, yon fwa yo ekri yon kle nan yon plas, li pa deplase ankò. Menm lè yon eleman efase nan tab la, kle a rete an plas, valè li tou senpleman vin vid. Sa vle di ke nou pa bezwen sèvi ak yon operasyon ekri atomik pou valè plas la, paske li pa enpòtan si valè aktyèl la vid oswa ou pa - li ap toujou vin vid.
Redimansyonman yon tab hash
Ou ka chanje gwosè yon tab hash pa kreye yon tab pi gwo epi mete eleman ki pa vid nan ansyen tab la nan li. Mwen pa t aplike fonksyonalite sa a paske mwen te vle kenbe kòd echantiyon an senp. Anplis, nan pwogram CUDA, alokasyon memwa souvan fèt nan kòd lame a olye ke nan nwayo CUDA a.
Atik la dekri kijan pou modifye yon estrikti done fèmen ak fèmen.
Konpetitivite
Nan fragments kòd fonksyon ki anwo yo gpu_hashtable_insert(), _lookup() и _delete() trete yon pè kle-valè alafwa. Ak pi ba gpu_hashtable_insert(), _lookup() и _delete() trete yon etalaj de pè an paralèl, chak pè nan yon fil egzekisyon GPU separe:
// 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);
}
}
Tablo hash ki reziste fèmen ak fèmen sipòte konkouran foure, rechèch, ak efase. Paske pè kle-valè yo toujou nan youn nan kat eta ak kle yo pa deplase, tab la garanti kòrèkteman menm lè diferan kalite operasyon yo itilize ansanm.
Sepandan, si nou trete yon pakèt ensèsyon ak efase an paralèl, epi si etalaj la antre nan pè gen kle kopi, Lè sa a, nou pa yo pral kapab predi ki pè yo pral "genyen" - yo pral ekri sou tab la hash dènye. Ann di nou rele kòd ensèsyon an ak yon etalaj opinyon pè A/0 B/1 A/2 C/3 A/4. Lè kòd la fini, pè B/1 и C/3 yo garanti yo dwe prezan nan tablo a, men an menm tan an nenpòt nan pè yo ap parèt nan li A/0, A/2 oswa A/4. Sa a ka oswa pa ka yon pwoblèm - li tout depann sou aplikasyon an. Ou ka konnen davans ke pa gen okenn kle kopi nan etalaj la antre, oswa ou ka pa pran swen ki valè yo te ekri dènye.
Si sa a se yon pwoblèm pou ou, Lè sa a, ou bezwen separe pè yo kopi nan diferan apèl sistèm CUDA. Nan CUDA, nenpòt operasyon ki rele nwayo a toujou fini anvan pwochen apèl nwayo a (omwen nan yon sèl fil. Nan diferan fil, nwayo yo egzekite an paralèl). Nan egzanp ki anwo a, si ou rele yon sèl nwayo ak A/0 B/1 A/2 C/3, ak lòt la avèk A/4, Lè sa a, kle a A pral jwenn valè a 4.
Koulye a, kite a pale sou si fonksyon yo ta dwe lookup() и delete() sèvi ak yon pwent senp oswa temèt nan yon etalaj de pè nan tablo hash la. Di ke:
Konpilatè a ka chwazi pou optimize lekti ak ekri nan memwa global oswa pataje... Optimizasyon sa yo ka enfim lè l sèvi avèk mo kle a.
volatile: ... nenpòt referans a varyab sa a konpile nan yon memwa reyèl lekti oswa ekri enstriksyon.
Konsiderasyon kòrèkteman pa mande pou aplikasyon an volatile. Si fil ekzekisyon an sèvi ak yon valè kachèt ki soti nan yon operasyon li pi bonè, Lè sa a, li pral sèvi ak enfòmasyon yon ti kras demode. Men, toujou, sa a se enfòmasyon ki soti nan eta ki kòrèk la nan tab la hash nan yon sèten moman nan apèl la nwayo. Si ou bezwen sèvi ak dènye enfòmasyon yo, ou ka itilize endèks la volatile, men Lè sa a, pèfòmans lan ap diminye yon ti kras: dapre tès mwen yo, lè efase 32 milyon dola eleman, vitès la diminye soti nan 500 milyon dola epresyon / sec a 450 milyon dola epresyon / sec.
Pwodiktivite
Nan tès la pou mete 64 milyon eleman ak efase 32 milyon nan yo, konpetisyon ant std::unordered_map epi pa gen pratikman pa gen okenn tab hash pou GPU la:
std::unordered_map te pase 70 ms mete ak retire eleman ak Lè sa a, libere yo unordered_map (Debarase m de dè milyon de eleman pran anpil tan, paske andedan unordered_map alokasyon memwa miltip fèt). Onètman pale, std:unordered_map restriksyon konplètman diferan. Li se yon sèl fil CPU nan ekzekisyon, sipòte kle-valè nan nenpòt ki gwosè, fè byen nan to itilizasyon segondè, epi li montre pèfòmans ki estab apre sipresyon miltip.
Dire tab la pou GPU ak kominikasyon ant pwogram yo te 984 ms. Sa a gen ladann tan ki pase mete tab la nan memwa epi efase li (atribye 1 GB memwa yon sèl fwa, ki pran kèk tan nan CUDA), mete ak efase eleman, ak iterasyon sou yo. Tout kopi pou ale ak pou soti nan memwa kat videyo yo tou pran an kont.
Tablo hash la tèt li te pran 271 ms pou konplete. Sa a gen ladan tan ki pase kat videyo a mete ak efase eleman, epi li pa pran an kont tan an te pase kopye nan memwa ak iterasyon sou tab la ki kapab lakòz. Si tab la GPU ap viv pou yon tan long, oswa si tab la hash ki genyen antyèman nan memwa a nan kat videyo a (pa egzanp, yo kreye yon tab hash ki pral itilize pa lòt kòd GPU epi yo pa processeur santral la), Lè sa a rezilta tès la enpòtan.
Tablo hash pou yon kat videyo demontre pèfòmans segondè akòz gwo debi ak paralelizasyon aktif.
Limit
Achitekti tab hash la gen kèk pwoblèm yo dwe okouran de:
- Sondaj lineyè antrave pa gwoupman, ki lakòz kle yo nan tab la yo dwe mete mwens pase parfe.
- Kle yo pa retire lè l sèvi avèk fonksyon an
deleteak sou tan yo ankonbre tab la.
Kòm yon rezilta, pèfòmans nan yon tab hash ka piti piti degrade, espesyalman si li egziste pou yon tan long epi li gen foure anpil ak efase. Youn nan fason yo bese dezavantaj sa yo se rehash nan yon nouvo tab ak yon to itilizasyon jistis ba ak filtre soti kle yo retire pandan rehashing la.
Pou ilistre pwoblèm ki dekri yo, mwen pral sèvi ak kòd ki anwo a pou kreye yon tab ki gen 128 milyon eleman ak bouk nan 4 milyon eleman jiskaske mwen ranpli 124 milyon fant (pousantaj itilizasyon apeprè 0,96). Isit la se tablo rezilta a, chak ranje se yon apèl nwayo CUDA pou mete 4 milyon nouvo eleman nan yon tab hash:
Pousantaj itilizasyon
Ensèsyon dire 4 eleman
0,00
11,608448 ms (361,314798 milyon kle/seg.)
0,03
11,751424 ms (356,918799 milyon kle/seg.)
0,06
11,942592 ms (351,205515 milyon kle/seg.)
0,09
12,081120 ms (347,178429 milyon kle/seg.)
0,12
12,242560 ms (342,600233 milyon kle/seg.)
0,16
12,396448 ms (338,347235 milyon kle/seg.)
0,19
12,533024 ms (334,660176 milyon kle/seg.)
0,22
12,703328 ms (330,173626 milyon kle/seg.)
0,25
12,884512 ms (325,530693 milyon kle/seg.)
0,28
13,033472 ms (321,810182 milyon kle/seg.)
0,31
13,239296 ms (316,807174 milyon kle/seg.)
0,34
13,392448 ms (313,184256 milyon kle/seg.)
0,37
13,624000 ms (307,861434 milyon kle/seg.)
0,41
13,875520 ms (302,280855 milyon kle/seg.)
0,44
14,126528 ms (296,909756 milyon kle/seg.)
0,47
14,399328 ms (291,284699 milyon kle/seg.)
0,50
14,690304 ms (285,515123 milyon kle/seg.)
0,53
15,039136 ms (278,892623 milyon kle/seg.)
0,56
15,478656 ms (270,973402 milyon kle/seg.)
0,59
15,985664 ms (262,379092 milyon kle/seg.)
0,62
16,668673 ms (251,627968 milyon kle/seg.)
0,66
17,587200 ms (238,486174 milyon kle/seg.)
0,69
18,690048 ms (224,413765 milyon kle/seg.)
0,72
20,278816 ms (206,831789 milyon kle/seg.)
0,75
22,545408 ms (186,038058 milyon kle/seg.)
0,78
26,053312 ms (160,989275 milyon kle/seg.)
0,81
31,895008 ms (131,503463 milyon kle/seg.)
0,84
42,103294 ms (99,619378 milyon kle/seg.)
0,87
61,849056 ms (67,815164 milyon kle/seg.)
0,90
105,695999 ms (39,682713 milyon kle/seg.)
0,94
240,204636 ms (17,461378 milyon kle/seg.)
Kòm itilizasyon ogmante, pèfòmans diminye. Sa a se pa dezirab nan pifò ka yo. Si yon aplikasyon mete eleman nan yon tab epi li jete yo (pa egzanp, lè w ap konte mo nan yon liv), sa a se pa yon pwoblèm. Men, si aplikasyon an sèvi ak yon tablo hash ki dire lontan (pa egzanp, nan yon editè grafik pou estoke pati ki pa vid nan imaj kote itilizatè a souvan foure ak efase enfòmasyon), Lè sa a, konpòtman sa a ka pwoblèm.
Ak mezire pwofondè sonde tab hash apre 64 milyon dola foure (faktè itilizasyon 0,5). Pwofondè an mwayèn te 0,4774, kidonk pifò kle yo te swa nan pi bon plas posib oswa yon sèl plas lwen pi bon pozisyon an. Pwofondè maksimòm son an te 60.
Lè sa a, mwen mezire pwofondè sonde sou yon tab ak 124 milyon foure (faktè itilizasyon 0,97). Pwofondè an mwayèn te deja 10,1757, ak maksimòm la - 6474 (!!). Pèfòmans deteksyon lineyè tonbe siyifikativman nan to itilizasyon segondè.
Li pi bon pou kenbe to itilizasyon tab hash sa a ba. Men, Lè sa a, nou ogmante pèfòmans nan depans lan nan konsomasyon memwa. Erezman, nan ka kle 32-bit ak valè, sa ka jistifye. Si nan egzanp ki anwo a, nan yon tablo ki gen 128 milyon eleman, nou kenbe faktè itilizasyon 0,25, Lè sa a, nou ka mete pa plis pase 32 milyon eleman nan li, epi rès 96 milyon fant yo pral pèdi - 8 bytes pou chak pè. , 768 MB nan memwa pèdi.
Tanpri sonje ke nou ap pale de pèt memwa kat videyo, ki se yon resous ki gen plis valè pase memwa sistèm. Malgre ke pifò kat grafik Desktop modèn ki sipòte CUDA gen omwen 4 GB memwa (nan moman sa a, NVIDIA 2080 Ti a gen 11 GB), li ta toujou pa desizyon ki pi saj pou pèdi kantite sa yo.
Apre sa, mwen pral ekri plis sou kreye tab hash pou kat videyo ki pa gen pwoblèm ak pwofondè sonde, osi byen ke fason yo reitilize fant efase.
Sonde pwofondè mezi
Pou detèmine pwofondè sonde yon kle, nou ka ekstrè hash kle a (endèks tab ideyal li yo) nan endèks tab aktyèl li yo:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Akòz majik la nan nimewo binè konpleman de de a ak lefèt ke kapasite nan tab la hash se de nan pouvwa a de, apwòch sa a ap travay menm lè endèks kle a deplase nan kòmansman an nan tab la. Ann pran yon kle ki hache a 1, men ki antre nan plas 3. Lè sa a, pou yon tab ki gen kapasite 4 nou jwenn (3 — 1) & 3, ki ekivalan a 2.
Konklizyon
Si w gen kesyon oswa kòmantè, tanpri imèl mwen nan oswa louvri yon nouvo sijè nan.
Kòd sa a te ekri sou enspirasyon nan atik ekselan:
Nan tan kap vini an, mwen pral kontinye ekri sou aplikasyon tab hash pou kat videyo ak analize pèfòmans yo. Plan mwen yo enkli chenn, Robin Hood hachaj, ak hachage koukou lè l sèvi avèk operasyon atomik nan estrikti done ki zanmitay GPU.
Sous: www.habr.com