Mi afiŝis ĝin sur Github .
Ĝi estas simpla GPU-haŝtabelo kapabla prilabori centojn da milionoj da enmetoj je sekundo. Sur mia tekkomputilo NVIDIA GTX 1060, la kodo enigas 64 milionojn hazarde generitaj ŝlosil-valoraj paroj en ĉirkaŭ 210 ms kaj forigas 32 milionojn da paroj en ĉirkaŭ 64 ms.
Tio estas, la rapideco sur tekkomputilo estas proksimume 300 milionoj da enigaĵoj/sek kaj 500 milionoj da forigoj/sec.
La tabelo estas skribita en CUDA, kvankam la sama tekniko povas esti aplikita al HLSL aŭ GLSL. La efektivigo havas plurajn limojn por certigi altan rendimenton sur vidkarto:
- Nur 32-bitaj ŝlosiloj kaj la samaj valoroj estas procesitaj.
- La hashtablo havas fiksan grandecon.
- Kaj ĉi tiu grandeco devas esti egala al du al la potenco.
Por ŝlosiloj kaj valoroj, vi devas rezervi simplan limmarkon (en la supra kodo ĉi tio estas 0xffffffff).
Hash-tablo sen seruroj
La hashtabelo uzas malferman adresadon kun , tio estas, ĝi estas simple tabelo de ŝlosil-valoraj paroj, kiu estas konservita en memoro kaj havas superan kaŝmemorefikecon. La samon ne povas diri pri ĉenado, kiu implicas serĉi montrilon en ligita listo. Hashtabelo estas simpla tabelo stokanta elementojn KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
La grandeco de la tablo estas potenco de du, ne unua nombro, ĉar unu rapida instrukcio sufiĉas por apliki la pow2/AND-maskon, sed la modula funkciigisto estas multe pli malrapida. Tio estas grava en la kazo de linia sondado, ĉar en linia tabelrigardo la fendeto-indekso devas esti envolvita en ĉiu fendeto. Kaj kiel rezulto, la kosto de la operacio estas aldonita modulo en ĉiu fendeto.
La tabelo nur konservas la ŝlosilon kaj valoron por ĉiu elemento, ne hash de la ŝlosilo. Ĉar la tabelo nur stokas 32-bitajn ŝlosilojn, la haŝo estas kalkulita tre rapide. La supra kodo uzas la Murmur3-haŝiŝon, kiu nur faras kelkajn ŝanĝojn, XOR-ojn kaj multiplikojn.
La hashtabelo uzas ŝlosajn protektoteknikojn kiuj estas sendependaj de memorordo. Eĉ se iuj skribaj operacioj interrompas la ordon de aliaj tiaj operacioj, la hashtabelo ankoraŭ konservos la ĝustan staton. Pri ĉi tio ni parolos sube. La tekniko funkcias bonege kun vidkartoj, kiuj kuras milojn da fadenoj samtempe.
La ŝlosiloj kaj valoroj en la hashtabelo estas pravigitaj por malplenigi.
La kodo povas esti modifita por manipuli ankaŭ 64-bitajn ŝlosilojn kaj valorojn. Ŝlosiloj postulas atomlegadon, skribon kaj kompar-kaj-interŝanĝajn operaciojn. Kaj valoroj postulas atomajn legajn kaj skribajn operaciojn. Feliĉe, en CUDA, legado-skriba operacioj por 32- kaj 64-bitaj valoroj estas atomaj kondiĉe ke ili estas nature vicigitaj (vidu sube). ), kaj modernaj vidkartoj subtenas 64-bitajn atomajn kompar-kaj-interŝanĝajn operaciojn. Kompreneble, kiam moviĝas al 64 bitoj, rendimento iomete malpliiĝos.
Hash-tabelo stato
Ĉiu ŝlosil-valora paro en hashtabelo povas havi unu el kvar statoj:
- Ŝlosilo kaj valoro estas malplenaj. En ĉi tiu stato, la hashtabelo estas pravigita.
- La ŝlosilo estas notita, sed la valoro ankoraŭ ne estas skribita. Se alia fadeno nuntempe legas datumojn, ĝi tiam revenas malplena. Ĉi tio estas normala, la sama afero estus okazinta se alia fadeno de ekzekuto funkcius iom pli frue, kaj ni parolas pri samtempa datumstrukturo.
- Kaj la ŝlosilo kaj la valoro estas registritaj.
- La valoro disponeblas por aliaj fadenoj de ekzekuto, sed la ŝlosilo ankoraŭ ne estas. Tio povas okazi ĉar la CUDA-programa modelo havas loze ordonitan memormodelon. Ĉi tio estas normala; ĉiuokaze, la ŝlosilo ankoraŭ estas malplena, eĉ se la valoro ne plu estas tia.
Grava nuanco estas, ke post kiam la ŝlosilo estas skribita al la fendo, ĝi ne plu moviĝas - eĉ se la ŝlosilo estas forigita, ni parolos pri tio ĉi sube.
La hashtabelkodo eĉ funkcias kun loze ordigitaj memormodeloj en kiuj la ordo en kiu memoro estas legita kaj skribita estas nekonata. Dum ni rigardas enmeton, serĉon kaj forigon en hashtabelo, memoru, ke ĉiu ŝlosil-valora paro estas en unu el la kvar statoj priskribitaj supre.
Enmetado en hashtabelon
La CUDA-funkcio, kiu enmetas ŝlosil-valorajn parojn en hashtabelon, aspektas jene:
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);
}
}
Por enmeti ŝlosilon, la kodo ripetas tra la hashtabelo komencante per la hash de la enigita ŝlosilo. Ĉiu fendeto en la tabelo elfaras atoman kompar-kaj-interŝanĝan operacion kiu komparas la ŝlosilon en tiu fendo al malplena. Se malkongruo estas detektita, la ŝlosilo en la fendo estas ĝisdatigita kun la enigita ŝlosilo, kaj tiam la origina fendoŝlosilo estas resendita. Se ĉi tiu origina ŝlosilo estis malplena aŭ kongruis kun la enigita ŝlosilo, tiam la kodo trovis taŭgan fendon por enmeto kaj enigis la enigitan valoron en la fendo.
Se en unu kernovoko gpu_hashtable_insert() estas pluraj elementoj kun la sama ŝlosilo, tiam iu ajn el iliaj valoroj povas esti skribita al la ŝlosila fendo. Ĉi tio estas konsiderata normala: unu el la ŝlosilvaloraj skriboj dum la voko sukcesos, sed ĉar ĉio ĉi okazas paralele ene de pluraj fadenoj de ekzekuto, ni ne povas antaŭdiri, kiu memorskribo estos la lasta.
Hash-tabelserĉo
Kodo por serĉi ŝlosilojn:
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);
}
}
Por trovi la valoron de ŝlosilo stokita en tabelo, ni ripetas tra la tabelo komencante per la hash de la ŝlosilo, kiun ni serĉas. En ĉiu fendo, ni kontrolas ĉu la ŝlosilo estas tiu, kiun ni serĉas, kaj se jes, ni resendas ĝian valoron. Ni ankaŭ kontrolas ĉu la ŝlosilo estas malplena, kaj se jes, ni ĉesas la serĉon.
Se ni ne povas trovi la ŝlosilon, la kodo resendas malplenan valoron.
Ĉiuj ĉi tiuj serĉaj operacioj povas esti faritaj samtempe per enmetoj kaj forigoj. Ĉiu paro en la tabelo havos unu el la kvar statoj priskribitaj supre por la fluo.
Forigo en hashtabelo
Kodo por forigi ŝlosilojn:
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);
}
}
Forigo de ŝlosilo estas farita en nekutima maniero: ni lasas la ŝlosilon en la tabelo kaj markas ĝian valoron (ne la ŝlosilon mem) kiel malplena. Ĉi tiu kodo estas tre simila al lookup(), krom ke kiam kongruo estas trovita sur ŝlosilo, ĝi faras ĝian valoron malplena.
Kiel menciite supre, post kiam ŝlosilo estas skribita al fendo, ĝi ne plu estas movita. Eĉ kiam elemento estas forigita de la tabelo, la ŝlosilo restas en loko, ĝia valoro simple fariĝas malplena. Tio signifas, ke ni ne bezonas uzi atoman skriban operacion por la fendovaloro, ĉar ne gravas ĉu la nuna valoro estas malplena aŭ ne - ĝi ankoraŭ malpleniĝos.
Regrandigo de hashtabelo
Vi povas ŝanĝi la grandecon de hashtabelo kreante pli grandan tablon kaj enmetante nemalplenajn elementojn de la malnova tablo en ĝin. Mi ne efektivigis ĉi tiun funkcion ĉar mi volis konservi la ekzemplan kodon simpla. Krome, en CUDA-programoj, memorasigno ofte estas farita en la gastiga kodo prefere ol en la CUDA-kerno.
En la artikolo priskribas kiel modifi tian serurprotektan datumstrukturon.
Konkuremo
En ĉi-supraj funkcio-kodfragmentoj gpu_hashtable_insert(), _lookup() и _delete() procesi unu ŝlosil-valoran paron samtempe. Kaj pli malalta gpu_hashtable_insert(), _lookup() и _delete() prilaboru aron da paroj paralele, ĉiu paro en aparta GPU-ekzekuta fadeno:
// 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);
}
}
La serurrezista hashtabelo subtenas samtempajn enmetojn, serĉojn kaj forigojn. Ĉar ŝlosil-valoraj paroj ĉiam estas en unu el kvar statoj kaj la ŝlosiloj ne moviĝas, la tabelo garantias ĝustecon eĉ kiam malsamaj specoj de operacioj estas uzataj samtempe.
Tamen, se ni procesas aron da enmetoj kaj forigoj paralele, kaj se la eniga tabelo de paroj enhavas duplikatajn ŝlosilojn, tiam ni ne povos antaŭdiri, kiuj paroj "venkos" — estos skribitaj laste al la hashtabelo. Ni diru, ke ni vokis la enigkodon kun eniga tabelo de paroj A/0 B/1 A/2 C/3 A/4. Kiam la kodo finiĝas, pariĝas B/1 и C/3 estas garantiitaj por ĉeesti en la tabelo, sed samtempe iu ajn el la paroj aperos en ĝi A/0, A/2 aŭ A/4. Ĉi tio eble aŭ ne estas problemo - ĉio dependas de la aplikaĵo. Vi eble scias anticipe, ke ne estas duplikataj ŝlosiloj en la eniga tabelo, aŭ eble vi ne zorgas pri kiu valoro estis skribita laste.
Se ĉi tio estas problemo por vi, tiam vi devas apartigi la duplikatajn parojn en malsamajn CUDA-sistemvokojn. En CUDA, ĉiu operacio kiu vokas la kernon ĉiam finiĝas antaŭ la sekva kernovoko (almenaŭ ene de unu fadeno. En malsamaj fadenoj, kernoj estas ekzekutitaj paralele). En la supra ekzemplo, se vi nomas unu kernon per A/0 B/1 A/2 C/3, kaj la alia kun A/4, tiam la ŝlosilo A ricevos la valoron 4.
Nun ni parolu pri ĉu funkcioj devus lookup() и delete() uzu simplan aŭ volatilan montrilon al tabelo de paroj en la hashtabelo. deklaras ke:
La kompililo povas elekti optimumigi legadojn kaj skribojn al tutmonda aŭ komuna memoro... Ĉi tiuj optimumigoj povas esti malŝaltitaj uzante la ŝlosilvorton
volatile: ... ajna referenco al ĉi tiu variablo estas kompilita en realan memoran instrukcion de legado aŭ skribado.
Ĝustecaj konsideroj ne postulas aplikon volatile. Se la ekzekutfadeno uzas kaŝmemorigitan valoron de pli frua legita operacio, tiam ĝi uzos iomete malnoviĝintajn informojn. Sed tamen, ĉi tio estas informo de la ĝusta stato de la hashtabelo en certa momento de la kernovoko. Se vi bezonas uzi la plej novajn informojn, vi povas uzi la indekson volatile, sed tiam la rendimento iomete malpliiĝos: laŭ miaj provoj, kiam vi forigas 32 milionojn da elementoj, la rapideco malpliiĝis de 500 milionoj da forigoj/sek al 450 milionoj da forigoj/sek.
Produkteco
En la provo por enmeti 64 milionojn da elementoj kaj forigi 32 milionojn da ili, konkurenco inter std::unordered_map kaj preskaŭ ne ekzistas hashtabelo por la GPU:
std::unordered_map pasigis 70 691 ms enmeti kaj forigi elementojn kaj poste liberigi ilin unordered_map (forigi milionojn da elementoj bezonas multan tempon, ĉar interne unordered_map multoblaj memor-asignoj estas faritaj). Sincere parolante, std:unordered_map tute malsamaj limigoj. Ĝi estas unuopa CPU-fadeno de ekzekuto, subtenas ŝlosilvalorojn de ajna grandeco, funkcias bone ĉe altaj utiligprocentoj kaj montras stabilan agadon post multoblaj forigoj.
La daŭro de la hashtabelo por la GPU kaj inter-programa komunikado estis 984 ms. Ĉi tio inkluzivas la tempon pasigitan metante la tablon en memoron kaj forigante ĝin (asignante 1 GB da memoro unufoje, kio prenas iom da tempo en CUDA), enigante kaj forigante elementojn, kaj ripetante super ili. Ĉiuj kopioj al kaj de la memoro de videokarto ankaŭ estas enkalkulitaj.
La hashtabelo mem daŭris 271 ms por kompletigi. Ĉi tio inkluzivas la tempon pasigitan de la vidkarto enigante kaj forigante elementojn, kaj ne konsideras la tempon pasigitan kopiante en memoron kaj ripetanta super la rezulta tablo. Se la GPU-tabelo vivas longe, aŭ se la hash-tabelo estas tute enhavita en la memoro de la vidkarto (ekzemple, por krei hash-tabelon, kiu estos uzata de alia GPU-kodo kaj ne de la centra procesoro), tiam la testrezulto estas grava.
La hashtabelo por vidkarto montras altan rendimenton pro alta trafluo kaj aktiva paraleligo.
mankoj
La arkitekturo de hashtabelo havas kelkajn problemojn por esti konscia pri:
- Lineara sondado estas malhelpita per amasigado, kiu igas la ŝlosilojn en la tabelon esti metitaj malpli ol perfekte.
- Ŝlosiloj ne estas forigitaj per la funkcio
deletekaj kun la tempo ili malordigas la tablon.
Kiel rezulto, la agado de hashtablo povas iom post iom malboniĝi, precipe se ĝi ekzistas dum longa tempo kaj havas multajn enmetojn kaj forigojn. Unu maniero mildigi ĉi tiujn malavantaĝojn estas rehash en novan tabelon kun sufiĉe malalta utiliga indico kaj filtri la forigitajn ŝlosilojn dum la rehashing.
Por ilustri la priskribitajn aferojn, mi uzos la ĉi-supran kodon por krei tabelon kun 128 milionoj da elementoj kaj trapasi 4 milionojn da elementoj ĝis mi plenigos 124 milionojn da fendoj (uzoprocento de ĉirkaŭ 0,96). Jen la rezulttabelo, ĉiu vico estas CUDA-kernovoko por enmeti 4 milionojn da novaj elementoj en unu hashtabelo:
Indice de uzado
Enmeto daŭro 4 elementoj
0,00
11,608448 ms (361,314798 milionoj da ŝlosiloj/sek.)
0,03
11,751424 ms (356,918799 milionoj da ŝlosiloj/sek.)
0,06
11,942592 ms (351,205515 milionoj da ŝlosiloj/sek.)
0,09
12,081120 ms (347,178429 milionoj da ŝlosiloj/sek.)
0,12
12,242560 ms (342,600233 milionoj da ŝlosiloj/sek.)
0,16
12,396448 ms (338,347235 milionoj da ŝlosiloj/sek.)
0,19
12,533024 ms (334,660176 milionoj da ŝlosiloj/sek.)
0,22
12,703328 ms (330,173626 milionoj da ŝlosiloj/sek.)
0,25
12,884512 ms (325,530693 milionoj da ŝlosiloj/sek.)
0,28
13,033472 ms (321,810182 milionoj da ŝlosiloj/sek.)
0,31
13,239296 ms (316,807174 milionoj da ŝlosiloj/sek.)
0,34
13,392448 ms (313,184256 milionoj da ŝlosiloj/sek.)
0,37
13,624000 ms (307,861434 milionoj da ŝlosiloj/sek.)
0,41
13,875520 ms (302,280855 milionoj da ŝlosiloj/sek.)
0,44
14,126528 ms (296,909756 milionoj da ŝlosiloj/sek.)
0,47
14,399328 ms (291,284699 milionoj da ŝlosiloj/sek.)
0,50
14,690304 ms (285,515123 milionoj da ŝlosiloj/sek.)
0,53
15,039136 ms (278,892623 milionoj da ŝlosiloj/sek.)
0,56
15,478656 ms (270,973402 milionoj da ŝlosiloj/sek.)
0,59
15,985664 ms (262,379092 milionoj da ŝlosiloj/sek.)
0,62
16,668673 ms (251,627968 milionoj da ŝlosiloj/sek.)
0,66
17,587200 ms (238,486174 milionoj da ŝlosiloj/sek.)
0,69
18,690048 ms (224,413765 milionoj da ŝlosiloj/sek.)
0,72
20,278816 ms (206,831789 milionoj da ŝlosiloj/sek.)
0,75
22,545408 ms (186,038058 milionoj da ŝlosiloj/sek.)
0,78
26,053312 ms (160,989275 milionoj da ŝlosiloj/sek.)
0,81
31,895008 ms (131,503463 milionoj da ŝlosiloj/sek.)
0,84
42,103294 ms (99,619378 milionoj da ŝlosiloj/sek.)
0,87
61,849056 ms (67,815164 milionoj da ŝlosiloj/sek.)
0,90
105,695999 ms (39,682713 milionoj da ŝlosiloj/sek.)
0,94
240,204636 ms (17,461378 milionoj da ŝlosiloj/sek.)
Ĉar utiligo pliiĝas, rendimento malpliiĝas. Ĉi tio ne estas dezirinda en la plej multaj kazoj. Se aplikaĵo enmetas elementojn en tabelon kaj poste forĵetas ilin (ekzemple, kiam oni kalkulas vortojn en libro), tiam tio ne estas problemo. Sed se la aplikaĵo uzas longdaŭran hashtabelon (ekzemple, en grafika redaktilo por stoki nemalplenajn partojn de bildoj, kie la uzanto ofte enmetas kaj forigas informojn), tiam ĉi tiu konduto povas esti problema.
Kaj mezuris la profundon de sondado de hashtabelo post 64 milionoj da enmetoj (utiliga faktoro 0,5). La averaĝa profundo estis 0,4774, do plej multaj ŝlosiloj estis aŭ en la plej bona ebla fendo aŭ unu fendo for de la plej bona pozicio. La maksimuma sonadprofundo estis 60.
Mi tiam mezuris la sondan profundon sur tablo kun 124 milionoj da enigaĵoj (utiliga faktoro 0,97). La meza profundo jam estis 10,1757, kaj la maksimumo - 6474 (!!). Lineara senta efikeco malpliiĝas signife ĉe altaj utiligprocentoj.
Plej bone estas teni malaltan la utiligan indicon de ĉi tiu hashtabelo. Sed tiam ni pliigas rendimenton koste de memorkonsumo. Feliĉe, en la kazo de 32-bitaj ŝlosiloj kaj valoroj, tio povas esti pravigita. Se en la supra ekzemplo, en tabelo kun 128 milionoj da elementoj, ni konservas la uzfaktoron de 0,25, tiam ni povas meti ne pli ol 32 milionojn da elementoj en ĝi, kaj la ceteraj 96 milionoj da fendoj perdiĝos - 8 bajtoj por ĉiu paro. , 768 MB de perdita memoro.
Bonvolu noti, ke ni parolas pri la perdo de memoro de videokarto, kiu estas pli valora rimedo ol sistema memoro. Kvankam la plej multaj modernaj labortablaj grafikaj kartoj, kiuj subtenas CUDA, havas almenaŭ 4 GB da memoro (en la momento de skribado, la NVIDIA 2080 Ti havas 11 GB), ĝi ankoraŭ ne estus la plej saĝa decido perdi tiajn kvantojn.
Poste mi skribos pli pri kreado de haŝtabloj por vidkartoj, kiuj ne havas problemojn pri sondado de profundo, kaj ankaŭ manierojn reuzi forigitajn fendojn.
Sona profundmezuro
Por determini la sondan profundon de ŝlosilo, ni povas ĉerpi la haŝon de la ŝlosilo (ĝia ideala tabelindekso) de ĝia fakta tabelindekso:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Pro la magio de du-komplementaj binaraj nombroj kaj la fakto ke la kapablo de la hashtabelo estas du al la potenco de du, ĉi tiu aliro funkcios eĉ kiam la ŝlosila indekso estas movita al la komenco de la tablo. Ni prenu ŝlosilon kiu haŝis al 1, sed estas enigita en fendo 3. Tiam por tablo kun kapacito 4 ni ricevas (3 — 1) & 3, kiu egalas al 2.
konkludo
Se vi havas demandojn aŭ komentojn, bonvolu retpoŝti al mi ĉe aŭ malfermu novan temon en.
Ĉi tiu kodo estis skribita sub inspiro de bonegaj artikoloj:
Estonte, mi daŭre skribos pri realigoj de hashtabeloj por vidkartoj kaj analizos ilian agadon. Miaj planoj inkluzivas ĉenadon, Robin Hood-hakadon kaj kukolo-hakadon uzante atomajn operaciojn en datumstrukturoj kiuj estas amika GPU.
fonto: www.habr.com