Es ievietoju GitHub .
Å Ä« ir vienkÄrÅ”a hash tabula GPU, kas spÄj apstrÄdÄt simtiem miljonu ievietojumu sekundÄ. ManÄ klÄpjdatorÄ ar NVIDIA GTX 1060 kods ievieto 64 miljonus nejauÅ”i Ä£enerÄtu atslÄgas-vÄrtÄ«bas pÄru aptuveni 210 ms laikÄ un izdzÄÅ” 32 miljonus pÄru aptuveni 64 ms laikÄ.
Tas nozÄ«mÄ, ka Ätrums manÄ klÄpjdatorÄ ir aptuveni 300 milj. ievietojumu/sekunžu un 500 milj. dzÄÅ”anu/sekunžu.
Tabula ir uzrakstÄ«ta CUDA, lai gan to paÅ”u metodi var piemÄrot HLSL vai GLSL. Izpildei ir vairÄki ierobežojumi, kas nodroÅ”ina augstu veiktspÄju uz grafiskÄs kartes:
- Tiek apstrÄdÄtas tikai 32 bitu atslÄgas un tÄdas paÅ”as vÄrtÄ«bas.
- Hash tabulai ir fiksÄts lielums.
- Un Å”im izmÄram jÄbÅ«t vienÄdam ar diviem pakÄpÄ.
AtslÄgÄm un vÄrtÄ«bÄm ir jÄpiedÄvÄ vienkÄrÅ”s robežzÄ«mes marÄ·ieris (dotajÄ kodÄ tas ir 0xffffffff).
Hash tabula bez bloÄ·ÄÅ”anas
Hash tabulÄ tiek izmantota atvÄrtÄ adresÄcija ar , tas ir, tas ir vienkÄrÅ”i atslÄgas-vÄrtÄ«bas pÄru masÄ«vs, kas tiek glabÄts atmiÅÄ un tam ir lieliska keÅ”atmiÅas veiktspÄja. To nevar teikt par saistÄ«Å”anu (chaining), kas nozÄ«mÄ, ka jÄmeklÄ rÄdÄ«tÄjs saistÄ«tajÄ sarakstÄ. Hash tabula ir vienkÄrÅ”s masÄ«vs, kas glabÄ elementus. KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Tabulas izmÄram jÄbÅ«t divciparu skaitlim, nevis pirmskaitlim, jo, lai izmantotu pow2/AND masku, ir nepiecieÅ”ama tikai viena Ätra instrukcija, turpretÄ« modulÄrais operators strÄdÄ daudz lÄnÄk. Tas ir svarÄ«gi lineÄrÄs zondÄÅ”anas gadÄ«jumÄ, jo, veicot lineÄro meklÄÅ”anu tabulÄ, slotu indeksam katrÄ slotÄ jÄpiedalÄs. Un tÄ rezultÄtÄ katrÄ slotÄ pievienojas izmaksas par operÄcijas moduli.
Tabula glabÄ tikai atslÄgu un vÄrtÄ«bu katram elementam, nevis atslÄgas hash. TÄ kÄ tabula glabÄ tikai 32 bitu atslÄgas, hash tiek aprÄÄ·inÄts ļoti Ätri. DotajÄ kodÄ tiek izmantots Murmur3 hash, kas veic tikai dažus bitu slÄgumus, XOR un reizinÄÅ”anas operÄcijas.
Hashtabula izmanto metodi, lai aizsargÄtu pret bloÄ·ÄÅ”anu, kas nav atkarÄ«ga no datu izvietoÅ”anas secÄ«bas atmiÅÄ. Pat ja daži rakstÄ«Å”anas darbi pÄrkÄpj secÄ«bu citu Å”Ädu operÄciju attiecÄ«bÄ, hashtabula joprojÄm saglabÄs pareizu stÄvokli. Par to mÄs runÄsim vÄlÄk. Å Ä« metode lieliski darbojas videokartÄs, kur tÅ«kstoÅ”iem pavedienu darbojas konkurÄjoÅ”i.
AtslÄgas un vÄrtÄ«bas hashtabulÄ tiek inicializÄtas kÄ tukÅ”as.
Kodu var modificÄt, lai tas varÄtu apstrÄdÄt arÄ« 64 bitu atslÄgas un vÄrtÄ«bas. AtslÄgÄm ir nepiecieÅ”amas atomÄrÄs operÄcijas lasÄ«Å”anai, rakstÄ«Å”anai un salÄ«dzinÄÅ”anai ar apmaiÅu (compare-and-swap). Bet vÄrtÄ«bÄm nepiecieÅ”amas atomÄrÄs operÄcijas lasÄ«Å”anai un rakstÄ«Å”anai. Par laimi, CUDA lasÄ«Å”anas un rakstÄ«Å”anas operÄcijas 32 un 64 bitu vÄrtÄ«bÄm ir atomÄras, kamÄr tÄs ir dabiski izkÄrtotas (skat. ), un mÅ«sdienu videokartes atbalsta 64 bitu atomÄrÄs salÄ«dzinÄÅ”anas ar apmaiÅu operÄcijas. Protams, pÄrejot uz 64 biti, veiktspÄja nedaudz samazinÄs.
Hashtabulas stÄvoklis
Katram atslÄgas-un-vÄrtÄ«bas pÄrim hashtabulÄ var bÅ«t viens no Äetriem stÄvokļiem:
- AtslÄga un vÄrtÄ«ba ir tukÅ”as. Å ÄdÄ stÄvoklÄ« hashtabula tiek inicializÄta.
- AtslÄga ir rakstÄ«ta, bet vÄrtÄ«ba vÄl nav. Ja Å”ajÄ brÄ«dÄ« cits izpildes pavedienu nolasÄ«s datus, tad tas atgriezÄ«s tukÅ”u vÄrtÄ«bu. Tas ir normÄli, tÄpat notiktu, ja cits izpildes pavedienu bÅ«tu izpildÄ«jies nedaudz agrÄk, un mÄs runÄjam par konkurÄjoÅ”u datu struktÅ«ru.
- Gan atslÄga, gan vÄrtÄ«ba ir rakstÄ«tas.
- VÄrtÄ«ba ir pieejama citiem izpildes pavedieniem, bet atslÄga vÄl nav. Tas var notikt, jo programmatÅ«ras modelis CUDA paredz vÄji sakÄrtotu atmiÅas modeli. Tas ir normÄli, jebkurÄ gadÄ«jumÄ atslÄga joprojÄm ir tukÅ”a pat tad, ja vÄrtÄ«ba tÄda jau nav.
SvarÄ«gs nianses ir tas, ka tiklÄ«dz atslÄga ir rakstÄ«ta slotÄ, tÄ vairs neatkÄpjas ā pat ja atslÄga tiks dzÄsta, mÄs par to runÄsim vÄlÄk.
Hashtabulas kods darbojas pat ar vÄji sakÄrtotiem atmiÅas modeļiem, kuros nav zinÄma lasÄ«Å”anas un rakstÄ«Å”anas secÄ«ba. Kad mÄs apspriedÄ«sim ievadÄ«Å”anu, meklÄÅ”anu un dzÄÅ”anu hashtabulÄ, atcerieties, ka katrs atslÄgas-un-vÄrtÄ«bas pÄris atrodas vienÄ no iepriekÅ” minÄtajiem Äetriem stÄvokļiem.
IevadÄ«Å”ana hashtabulÄ
CUDA funkcija, kas ievieto atslÄgas-vÄrtÄ«bu pÄrus hash tabulÄ, izskatÄs Å”Ädi:
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);
}
}
Lai ievietotu atslÄgu, kods iterÄ cauri hash tabulas masÄ«vam, sÄkot no ievietotÄs atslÄgas hasha. Katram masÄ«va slotam tiek veikta atomÄra salÄ«dzinÄjuma ar apmaiÅu operÄcija, kurÄ atslÄga Å”ajÄ slotÄ tiek salÄ«dzinÄta ar tukÅ”u. Ja atrodam atŔķirÄ«bu, atslÄga slotÄ tiek atjauninÄta uz ievietojamo atslÄgu, un pÄc tam tiek atgriezta sÄkotnÄjÄ atslÄga no slotÄ. Ja Ŕī oriÄ£inÄlÄ atslÄga bija tukÅ”a vai atbilst ievietotajai atslÄgai, tas nozÄ«mÄ, ka kods ir atradis piemÄrotu slotu ievietoÅ”anai un ievieto slotÄ ievietoto vÄrtÄ«bu.
Ja vienÄ kodu izsaukumÄ gpu_hashtable_insert() ir vairÄki elementi ar to paÅ”u atslÄgu, tad jebkura no to vÄrtÄ«bÄm var tikt ierakstÄ«ta atslÄgas slotÄ. Tas tiek uzskatÄ«ts par normÄlu: viena no atslÄgas-vÄrtÄ«bas ierakstÄ«Å”anas operÄcijÄm izsaukuma laikÄ bÅ«s veiksmÄ«ga, bet, tÄ kÄ tas viss notiek paralÄli vairÄkos izpildes pavedienos, mÄs nevaram paredzÄt, kura rakstÄ«Å”anas operÄcija atmiÅÄ bÅ«s pÄdÄjÄ.
MeklÄÅ”ana hash tabulÄ
Kods atslÄgu meklÄÅ”anai:
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);
}
}
Lai atrastu atslÄgas vÄrtÄ«bu, kas glabÄjas tabulÄ, mÄs iterÄjam cauri masÄ«vam, sÄkot no meklÄtÄs atslÄgas hasha. Katram slotam mÄs pÄrbaudÄm, vai atslÄga ir tÄ, kuru meklÄjam, un, ja jÄ, atgriežam tÄs vÄrtÄ«bu. MÄs arÄ« pÄrbaudÄm, vai atslÄga ir tukÅ”a, un, ja jÄ, mÄs pÄrtraucam meklÄÅ”anu.
Ja mÄs nevaram atrast atslÄgu, tad kods atgriež tukÅ”u vÄrtÄ«bu.
Visas Ŕīs meklÄÅ”anas operÄcijas var tikt veiktas konkurÄjoÅ”i, ievietoÅ”anas un dzÄÅ”anas laikÄ. Katram pÄrim tabulÄ bÅ«s Äetri iepriekÅ” aprakstÄ«tie stÄvokļi konkrÄtam pavedienam.
DzÄÅ”ana hash tabulÄ
Kods atslÄgu dzÄÅ”anai:
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);
}
}
AtslÄgas dzÄÅ”ana notiek neparasti: mÄs atstÄjam atslÄgu tabulÄ un atzÄ«mÄjam tÄs vÄrtÄ«bu (ne paÅ”u atslÄgu) kÄ tukÅ”u. Å is kods ir ļoti lÄ«dzÄ«gs lookup(), izÅemot to, ka, atrodot atbilstÄ«bu atslÄgai, tÄ vÄrtÄ«ba kļūst par tukÅ”u.
KÄ minÄts iepriekÅ”, tiklÄ«dz atslÄga ir ierakstÄ«ta slotÄ, tÄ vairs nepÄrvietojas. Pat tad, ja elements tiek dzÄsts no tabulas, atslÄga paliek savÄ vietÄ, vienkÄrÅ”i tÄs vÄrtÄ«ba kļūst tukÅ”a. Tas nozÄ«mÄ, ka mums nav nepiecieÅ”ams izmantot atomÄro rakstÄ«Å”anas operÄciju slotÄ, jo nav svarÄ«gi, vai paÅ”reizÄjÄ vÄrtÄ«ba ir tukÅ”a vai nÄ ā tÄ tomÄr kļūs tukÅ”a.
Hesh-tabulas izmÄra maiÅa
MainÄ«t hesh-tabulas izmÄru var, izveidojot lielÄku tabulu un ievietojot tajÄ nepilnÄ«gÄs elementus no vecÄs tabulas. Es Å”o funkcionalitÄti neesmu Ä«stenojis, jo vÄlÄjos saglabÄt koda paraugu vienkÄrÅ”u. TurklÄt CUDA programmÄs atmiÅas pieŔķirÅ”ana bieži tiek veikta hosta kodÄ, nevis CUDA kodÄ.
RakstÄ apraksta, kÄ mainÄ«t Å”Ädu datu struktÅ«ru, kas ir aizsargÄta no bloÄ·ÄÅ”anÄs.
Konkurence
Ievietotajos iepriekÅ”ÄjÄ koda fragmentos funkcijas gpu_hashtable_insert(), _lookup() un _delete() apstrÄdÄ pa vienai atslÄgas-vÄrtÄ«bas pÄrim reizÄ. Bet zemÄk gpu_hashtable_insert(), _lookup() un _delete() apstrÄdÄ pÄru masÄ«vu paralÄli, katru pÄri atseviÅ”Ä·Ä GPU izpildes plÅ«smÄ:
// 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);
}
}
Hesh-tabula, kas ir aizsargÄta no bloÄ·ÄÅ”anÄs, atbalsta konkurÄtspÄjÄ«gas ievietoÅ”anas, meklÄÅ”anas un dzÄÅ”anas operÄcijas. TÄ kÄ atslÄgas-vÄrtÄ«bas pÄri vienmÄr atrodas vienÄ no Äetriem stÄvokļiem, un atslÄgas nepÄrvietojas, tabula garantÄ precizitÄti pat tad, ja tiek vienlaicÄ«gi lietotas dažÄdu veidu operÄcijas.
TomÄr, ja mÄs paralÄli apstrÄdÄjam ievietoÅ”anas un dzÄÅ”anas partiju, un ja iebraukÅ”anas pÄru masÄ«vÄ ir atkÄrtojumi, mÄs nevarÄsim paredzÄt, kura pÄra "uzvarÄs" ā tÄs tiks ierakstÄ«tas hesh-tabulÄ pÄdÄjÄs. PieÅemsim, ka mÄs izsaucÄm ievietoÅ”anas kodu ar iebraukÅ”anas masÄ«vu pÄriem A/0 B/1 A/2 C/3 A/4. Kad kods pabeigs darbÄ«bu, pÄri B/1 un C/3 neapÅ”aubÄmi bÅ«s tabulÄ, bet tajÄ bÅ«s jebkura no pÄriem A/0, A/2 vai A/4Tas var bÅ«t problÄma vai arÄ« ne - viss atkarÄ«gs no pielietoÅ”anas. JÅ«s varat iepriekÅ” zinÄt, ka ievades masÄ«vÄ nav atkÄrtotu atslÄgu, vai arÄ« jums var bÅ«t vienalga, kura vÄrtÄ«ba ir ierakstÄ«ta pÄdÄjÄ.
Ja tas jums ir problÄma, jums ir jÄnodala atkÄrtotÄs pÄris no dažÄdÄm sistÄmas CUDA izsaukumiem. CUDA jebkura operÄcija ar kodolu izsaukumu vienmÄr beidzas pirms nÄkamÄ kodola izsaukuma (vismaz viena plÅ«smas ietvaros. AtŔķirÄ«gÄs plÅ«smÄs kodoli tiek izpildÄ«ti paralÄli). Ja iepriekÅ” minÄtajÄ piemÄrÄ izsaucat vienu kodolu ar A/0 B/1 A/2 C/3, un citu ar A/4, tad atslÄga A saÅems vÄrtÄ«bu 4.
Tagad parunÄsim par to, vai funkcijÄm lookup() un delete() jÄizmanto vienkÄrÅ”s (plain) vai mainÄ«gs (volatile) rÄdÄ«tÄjs uz pÄru masÄ«vu haÅ”a tabulÄ. apgalvo, ka:
Kompilatoram ir tiesÄ«bas optimizÄt lasÄ«Å”anas un rakstÄ«Å”anas operÄcijas uz globÄlo vai kopÄ«go atmiÅuā¦ Å Ä«s optimizÄcijas var atspÄjot, izmantojot atslÄgas vÄrdu
volatile: ... jebkura atsauce uz Å”o mainÄ«go tiek kompilÄta kÄ faktiskÄ lasÄ«Å”anas vai rakstÄ«Å”anas instrukcija atmiÅÄ.
Korekcijas apsvÄrumi neprasa piemÄrot volatile. Ja izpildes plÅ«sma izmanto keÅ”otu vÄrtÄ«bu no iepriekÅ”ÄjÄs lasÄ«Å”anas operÄcijas, tas nozÄ«mÄ, ka tÄ izmantos nedaudz novecojuÅ”u informÄciju. Bet tomÄr tas ir informÄcija no korekta stÄvokļa haÅ”a tabulÄ noteiktÄ kodola izsaukuma brÄ«dÄ«. Ja jums ir nepiecieÅ”ams izmantot visjaunÄko informÄciju, varat izmantot rÄdÄ«tÄju volatile, bet tad nedaudz samazinÄsies veiktspÄja: pÄc maniem testiem, izÅemot 32 miljonus elementu, Ätrums samazinÄjÄs no 500 miljoniem dzÄÅ”anu/sek lÄ«dz 450 miljoniem dzÄÅ”anu/sek.
VeiktspÄja
IevietoÅ”anas testÄ 64 miljonu elementu un 32 miljonu no tiem dzÄÅ”anas sacensÄ«ba starp std::unordered_map un haÅ”a tabulu GPU faktiski nepastÄv:
std::unordered_map tika pavadÄ«ti 70 691 ms elementu ievietoÅ”anai un dzÄÅ”anai, pÄc tam atbrÄ«voÅ”anai unordered_map (atbrÄ«voÅ”ana no miljoniem elementu aizÅem daudz laika, jo iekÅ”pusÄ tiek veiktas daudzas atmiÅas pieŔķirÅ”anÄs). GodÄ«gi sakot, pie unordered_map std::unordered_map std:unordered_map pilnÄ«gi citi ierobežojumi. Tas ir viens CPU izpildplÅ«sma, tas atbalsta atslÄgas-vÄrtÄ«bas jebkura izmÄra, labi darbojas ar augstiem izmantoÅ”anas koeficientiem un parÄda stabilu veiktspÄju pÄc daudzkÄrtÄjÄm dzÄÅ”anÄm.
Hash tabulas darbÄ«bas ilgums GPU un starpprogrammu mijiedarbÄ«bÄ bija 984 ms. Tas ietver laiku, kas tika pavadÄ«ts tabulas izvietoÅ”anai atmiÅÄ un tÄs dzÄÅ”anai (vienreizÄju 1 GB atmiÅas pieŔķirÅ”anu, kas CUDA ilgst zinÄmu laiku), elementu ievietoÅ”anai un dzÄÅ”anai, kÄ arÄ« iterÄÅ”anai pa tiem. TÄpat Åemti vÄrÄ visi kopÄÅ”anas procesi uz un no videokartes atmiÅas.
PaÅ”as hash tabulas darbÄ«ba aizÅÄma 271 ms. Tas ietver laiku, ko videokarte iztÄrÄja elementu ievietoÅ”anai un dzÄÅ”anai, un neietver laiku kopÄÅ”anai uz atmiÅu un iterÄÅ”anai pa iegÅ«to tabulu. Ja GPU tabula dzÄ«vo ilgi, vai ja hash tabula ir pilnÄ«gi videokartes atmiÅÄ (piemÄram, hash tabulas izveidoÅ”anai, ko izmantos cits GPU kods, nevis centrÄlais procesors), tad testa rezultÄti ir nozÄ«mÄ«gi.
Hash tabula videokartei demonstrÄ augstu veiktspÄju, pateicoties lielai caurlaidspÄjai un aktÄ«vai paralÄlajai izpildei.
Trūkumi
Hash tabulas arhitektÅ«rai ir dažas problÄmas, par kurÄm jÄatceras:
- LÄ«nijas izpÄte traucÄ klasterizÄcija, dÄļ kuras atslÄgas tabulÄ tiek novietotas tÄlu neideÄli.
- AtslÄgas netiek dzÄstas ar funkciju
izdzÄstun laika gaitÄ aizpilda tabulu.
KÄ rezultÄtÄ hash tabulas veiktspÄja var pakÄpeniski samazinÄties, Ä«paÅ”i, ja tÄ pastÄv ilgu laiku un tajÄ tiek veikti daudzi ievietojumi un dzÄÅ”anas procesi. Viens no veidiem, kÄ mazinÄt Å”os trÅ«kumus, ir atkÄrtota kopÅ”ana uz jaunu tabulu ar pietiekami zemu izmantoÅ”anas koeficientu un dzÄsto atslÄgu filtrÄcija atkÄrtotÄ kopÅ”anÄ.
Lai ilustrÄtu iepriekÅ” minÄtÄs problÄmas, izmantoÅ”u iepriekÅ” minÄto kodu, lai izveidotu tabulu ar 128 miljoniem elementu, cikliski ievietoÅ”u 4 miljonus elementu, lÄ«dz aizpildÄ«Å”u 124 miljonus slotu (izmantoÅ”anas koeficients apmÄram 0,96). Å eit ir rezultÄtu tabula, katra rinda ir CUDA kodola izsaukums ar 4 miljonu jauno elementu ievietoÅ”anu vienÄ hash tabulÄ:
IzmantoŔanas koeficients
IevietoŔanas ilgums 4 194 304 elementiem
0,00
11,608448 ms (361,314798 miljoni atslÄgu/s)
0,03
11,751424 ms (356,918799 miljoni atslÄgu/s)
0,06
11,942592 ms (351,205515 miljoni atslÄgu/s)
0,09
12,081120 ms (347,178429 miljoni atslÄgu/s)
0,12
12,242560 ms (342,600233 miljoni atslÄgu/s)
0,16
12,396448 ms (338,347235 miljoni atslÄgu/s)
0,19
12,533024 ms (334,660176 miljoni atslÄgu/s)
0,22
12,703328 ms (330,173626 miljoni atslÄgu/s)
0,25
12,884512 ms (325,530693 miljoni atslÄgu/s)
0,28
13,033472 ms (321,810182 miljoni atslÄgu/s)
0,31
13,239296 ms (316,807174 miljoni atslÄgu/s)
0,34
13,392448 ms (313,184256 miljoni atslÄgu/s)
0,37
13,624000 ms (307,861434 miljoni atslÄgu/s)
0,41
13,875520 ms (302,280855 miljoni atslÄgu/s)
0,44
14,126528 ms (296,909756 miljoni atslÄgu/s)
0,47
14,399328 ms (291,284699 miljoni atslÄgu/s)
0,50
14,690304 ms (285,515123 miljoni atslÄgu/s)
0,53
15,039136 ms (278,892623 miljoni atslÄgu/s)
0,56
15,478656 ms (270,973402 miljoni atslÄgu/s)
0,59
15,985664 ms (262,379092 miljoni atslÄgu/s)
0,62
16,668673 ms (251,627968 miljoni atslÄgu/s)
0,66
17,587200 ms (238,486174 miljoni atslÄgu/s)
0,69
18,690048 ms (224,413765 miljoni atslÄgu/s)
0,72
20,278816 ms (206,831789 miljoni atslÄgu/s)
0,75
22,545408 ms (186,038058 miljoni atslÄgu/s)
0,78
26,053312 ms (160,989275 miljoni atslÄgu/s)
0,81
31,895008 ms (131,503463 miljoni atslÄgu/s)
0,84
42,103294 ms (99,619378 miljoni atslÄgu/s)
0,87
61,849056 ms (67,815164 miljoni atslÄgu/s)
0,90
105,695999 ms (39,682713 miljoni atslÄgu/s)
0,94
240,204636 ms (17,461378 miljoni atslÄgu/s)
LielÄkam izmantoÅ”anas koeficientam samazinÄs veiktspÄja. Tas lielÄkajÄ daÄ¼Ä gadÄ«jumu ir nevÄlams. Ja lietotne ievieto elementus tabulÄ un pÄc tam tos noraida (piemÄram, skaitot vÄrdus grÄmatas laikÄ), tas nav problÄma. Bet ja lietotne izmanto ilgstoÅ”u hash tabulu (piemÄram, grafiskajÄ redaktorÄ, lai uzglabÄtu neiztukÅ”otas attÄlu daļas, kad lietotÄjs bieži ievieto un dzÄÅ” informÄciju), tad Å”Äda uzvedÄ«ba var radÄ«t neÄrtÄ«bas.
Un es izmÄrÄ«ju hash tabulas zondes dziļumu pÄc 64 miljoniem ievietoÅ”anas (izmantoÅ”anas koeficients 0,5). VidÄjÄ dziļums bija 0,4774, tÄdÄļ lielÄkÄ daļa atslÄgu bija vai nu labÄkajos iespÄjamos slotos, vai vienÄ slotÄ no labÄkÄs pozÄ«cijas. MaksimÄlÄ zondes dziļums bija 60.
PÄc tam es izmÄrÄ«ju zondes dziļumu tabulÄ ar 124 miljoniem ievietoÅ”anas (izmantoÅ”anas koeficients 0,97). VidÄjÄ dziļums jau bija 10,1757, bet maksimÄlÄ ā 6474 (!!). LineÄrÄs zondÄÅ”anas veiktspÄja strauji samazinÄs pie augstiem izmantoÅ”anas koeficientiem.
VislabÄk ir saglabÄt Å”ai hash tabulai zemu izmantoÅ”anas koeficients. TomÄr tÄdÄjÄdi mÄs paaugstinÄm veiktspÄju, palielinot atmiÅas patÄriÅu. Par laimi, 32 bitu atslÄgu un vÄrtÄ«bu gadÄ«jumÄ, tas var bÅ«t pamatots. Ja iepriekÅ” minÄtajÄ piemÄrÄ tabulÄ ar 128 miljoniem elementu saglabÄsim izmantoÅ”anas koeficientu 0,25, tad mÄs varÄsim tajÄ novietot ne vairÄk kÄ 32 miljonus elementu, bet pÄrÄjie 96 miljoni vietu tiks zaudÄti ā pa 8 baitiem uz katru pÄri, 768 Mb zaudÄtÄs atmiÅas.
LÅ«dzu, Åemiet vÄrÄ, ka Å”eit runa ir par videokartes zaudÄto atmiÅu, kas ir vÄrtÄ«gÄks resurss nekÄ sistÄmas atmiÅa. Lai gan lielÄkajai daļai mÅ«sdienu galddatoru videokartÄm, kas atbalsta CUDA, ir vismaz 4 Gb atmiÅas (rakstÄ«Å”anas brÄ«dÄ« NVIDIA 2080 Ti ir 11 Gb), tomÄr zaudÄt Å”Ädus apjomus nebÅ«s prÄtÄ«gi.
VÄlÄk es sÄ«kÄk rakstÄ«Å”u par hash tabulu izveidi videokartÄm, kurÄm nav problÄmu ar probe dziļumu, kÄ arÄ« par veidiem, kÄ atkÄrtoti izmantot dzÄsto vietu.
Probe dziļuma mÄrÄ«Å”ana
Lai noteiktu atslÄgas probe dziļumu, mÄs varam iegÅ«t atslÄgas hash (tÄs ideÄlo indeksu tabulÄ) no tÄ faktiskÄ tabulas indeksa:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
DÄļ divu binÄro skaitļu burvÄ«bas papildinÄÅ”anas kodÄ un fakts, ka hash tabulas ietilpÄ«ba ir divi spÄkÄ, Å”is pieejas darbosies pat tad, ja atslÄgas indekss tiek pÄrvietots tabulas sÄkumÄ. Å
emsim atslÄgu, kas hashÄjas 1, bet ir ievietota vietÄ 3. TÄdÄjÄdi tabulai ar ietilpÄ«bu 4 mÄs iegÅ«sim (3 ā 1) & 3, kas ir ekvivalents 2.
SecinÄjums
Ja jums ir jautÄjumi vai komentÄri, rakstiet man uz vai atveriet jaunu tÄmu.
Å is kods ir rakstÄ«ts iedvesmojoties no brÄ«niŔķīgÄm rakstiem:
NÄkotnÄ es turpinÄÅ”u rakstÄ«t par hash tabulu realizÄcijÄm videokartÄm un analizÄÅ”u to veiktspÄju. ManÄ plÄnÄ ir Ä·Ädes saistÄ«Å”ana, Robina Huda hashÄÅ”ana un kukaiÅ”a hashÄÅ”ana, izmantojot atomiskÄs operÄcijas datu struktÅ«rÄs, kuras ir piemÄrotas videokartÄm.
Avots: habr.com
Russian
Italian
German
Polish
Estonian
Albanian
Bulgarian
Croatian
Czech
Danish
Dutch
English
French
Hungarian
Latvian
Portuguese
Romanian
Slovak
Slovenian
Spanish
Swedish