ProHoster > Blogs > Administrācija > VienkārÅ”a jaucēj tabula GPU

VienkārÅ”a jaucēj tabula GPU

VienkārÅ”a jaucēj tabula GPU
Es to ievietoju Github jauns projekts VienkārŔa GPU hash tabula.

Tā ir vienkārÅ”a GPU hash tabula, kas spēj apstrādāt simtiem miljonu ieliktņu sekundē. Manā NVIDIA GTX 1060 klēpjdatorā kods ievieto 64 miljonus nejauÅ”i Ä£enerētu atslēgu vērtÄ«bu pāru aptuveni 210 ms laikā un noņem 32 miljonus pāru aptuveni 64 ms laikā.

Tas nozÄ«mē, ka klēpjdatora ātrums ir aptuveni 300 miljoni ievietoÅ”anas/sek un 500 miljoni dzÄ“Å”anas/sek.

Tabula ir uzrakstÄ«ta CUDA, lai gan to paÅ”u tehniku ā€‹ā€‹var izmantot HLSL vai GLSL. IevieÅ”anai ir vairāki ierobežojumi, lai nodroÅ”inātu augstu veiktspēju videokartē:

  • Tiek apstrādātas tikai 32 bitu atslēgas un tās paÅ”as vērtÄ«bas.
  • Hash tabulai ir fiksēts izmērs.
  • Un Å”im izmēram jābÅ«t vienādam ar diviem lÄ«dz jaudai.

Atslēgām un vērtÄ«bām ir jārezervē vienkārÅ”s norobežotāja marÄ·ieris (iepriekÅ” minētajā kodā tas ir 0xffffffff).

Hash galds bez slēdzenēm

Jaucējtabulā tiek izmantota atvērtā adrese ar lineārā zondÄ“Å”ana, tas ir, tas ir vienkārÅ”i atslēgu un vērtÄ«bu pāru masÄ«vs, kas tiek glabāts atmiņā un kam ir izcila keÅ”atmiņas veiktspēja. To paÅ”u nevar teikt par ķēdes pievienoÅ”anu, kas ietver rādÄ«tāja meklÄ“Å”anu saistÄ«tajā sarakstā. Hash tabula ir vienkārÅ”s masÄ«vs, kurā tiek glabāti elementi KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Tabulas izmērs ir pakāpē divi, nevis pirmskaitlis, jo pietiek ar vienu ātru instrukciju, lai uzliktu masku pow2/AND, bet moduļa operators ir daudz lēnāks. Tas ir svarÄ«gi lineārās zondÄ“Å”anas gadÄ«jumā, jo lineārās tabulas uzmeklÄ“Å”anā slota indekss ir jāiekļauj katrā slotā. Un rezultātā operācijas izmaksas tiek pievienotas modulo katrā slotā.

Tabulā tiek saglabāta tikai katra elementa atslēga un vērtÄ«ba, nevis atslēgas jaucējkods. Tā kā tabulā tiek saglabātas tikai 32 bitu atslēgas, hash tiek aprēķināts ļoti ātri. IepriekÅ” minētais kods izmanto Murmur3 hash, kas veic tikai dažas maiņas, XOR un reizināŔanu.

Jaucēj tabulā tiek izmantotas bloÄ·Ä“Å”anas aizsardzÄ«bas metodes, kas nav atkarÄ«gas no atmiņas secÄ«bas. Pat ja dažas rakstÄ«Å”anas darbÄ«bas izjauc citu Ŕādu darbÄ«bu secÄ«bu, hash tabula joprojām saglabās pareizo stāvokli. Par to mēs runāsim tālāk. Å Ä« tehnika lieliski darbojas ar videokartēm, kurās vienlaikus darbojas tÅ«kstoÅ”iem pavedienu.

Atslēgas un vērtÄ«bas hash tabulā tiek inicializētas, lai tās bÅ«tu tukÅ”as.

Kodu var modificēt, lai apstrādātu arÄ« 64 bitu atslēgas un vērtÄ«bas. Atslēgām ir nepiecieÅ”amas atomu lasÄ«Å”anas, rakstÄ«Å”anas un salÄ«dzināŔanas un apmaiņas darbÄ«bas. Un vērtÄ«bām ir nepiecieÅ”amas atomu lasÄ«Å”anas un rakstÄ«Å”anas darbÄ«bas. Par laimi, CUDA lasÄ«Å”anas un rakstÄ«Å”anas darbÄ«bas 32 un 64 bitu vērtÄ«bām ir atomāras, ja vien tās ir dabiski izlÄ«dzinātas (skatiet tālāk). Å”eit), un modernās videokartes atbalsta 64 bitu atomu salÄ«dzināŔanas un apmaiņas darbÄ«bas. Protams, pārejot uz 64 bitiem, veiktspēja nedaudz samazināsies.

JaukŔanas tabulas stāvoklis

Katram atslēgu un vērtÄ«bu pārim jaukÅ”anas tabulā var bÅ«t viens no četriem stāvokļiem:

  • Atslēga un vērtÄ«ba ir tukÅ”a. Å ajā stāvoklÄ« hash tabula tiek inicializēta.
  • Atslēga ir pierakstÄ«ta, bet vērtÄ«ba vēl nav uzrakstÄ«ta. Ja cits pavediens paÅ”laik lasa datus, tas atgriežas tukÅ”s. Tas ir normāli, tas pats bÅ«tu noticis, ja cits izpildes pavediens bÅ«tu darbojies nedaudz agrāk, un mēs runājam par vienlaicÄ«gu datu struktÅ«ru.
  • Tiek ierakstÄ«ta gan atslēga, gan vērtÄ«ba.
  • VērtÄ«ba ir pieejama citiem izpildes pavedieniem, bet atslēga vēl nav pieejama. Tas var notikt tāpēc, ka CUDA programmÄ“Å”anas modelim ir brÄ«vi sakārtots atmiņas modelis. Tas ir normāli; jebkurā gadÄ«jumā atslēga joprojām ir tukÅ”a, pat ja vērtÄ«ba vairs nav tāda.

SvarÄ«ga nianse ir tā, ka, tiklÄ«dz atslēga ir ierakstÄ«ta slotā, tā vairs nekustas ā€“ pat ja atslēga tiek izdzēsta, par to mēs runāsim tālāk.

Hash tabulas kods darbojas pat ar brÄ«vi sakārtotiem atmiņas modeļiem, kuros atmiņas lasÄ«Å”anas un rakstÄ«Å”anas secÄ«ba nav zināma. AplÅ«kojot ievietoÅ”anu, uzmeklÄ“Å”anu un dzÄ“Å”anu hash tabulā, atcerieties, ka katrs atslēgas un vērtÄ«bas pāris atrodas vienā no četriem iepriekÅ” aprakstÄ«tajiem stāvokļiem.

IevietoŔana hash tabulā

CUDA funkcija, kas ievieto atslēgu-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 atkārtojas caur hash tabulas masÄ«vu, sākot ar ievietotās atslēgas jaucējkodu. Katrs masÄ«va slots veic atomu salÄ«dzināŔanas un maiņas darbÄ«bu, kas salÄ«dzina Å”ajā slotā esoÅ”o atslēgu ar tukÅ”o. Ja tiek konstatēta neatbilstÄ«ba, slotā esoŔā atslēga tiek atjaunināta ar ievietoto atslēgu, un pēc tam tiek atgriezta sākotnējā slota atslēga. Ja Ŕī sākotnējā atslēga bija tukÅ”a vai atbilda ievietotajai atslēgai, kods atrada ievietoÅ”anai piemērotu slotu un ievietoja ievietoto vērtÄ«bu slotā.

Ja vienā kodola izsaukumā gpu_hashtable_insert() ir vairāki elementi ar vienu un to paÅ”u atslēgu, tad jebkuru no to vērtÄ«bām var ierakstÄ«t atslēgas slotā. Tas tiek uzskatÄ«ts par normālu: viena no atslēgas vērtÄ«bas ierakstiem zvana laikā izdosies, taču, tā kā tas viss notiek paralēli vairākos izpildes pavedienos, mēs nevaram paredzēt, kura atmiņas ierakstÄ«Å”ana bÅ«s pēdējā.

Hash tabulas meklēŔana

Kods taustiņu 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 tabulā saglabātās atslēgas vērtÄ«bu, mēs atkārtojam masÄ«vu, sākot ar meklētās atslēgas jaucējkodu. Katrā slotā mēs pārbaudām, vai atslēga ir tā, ko meklējam, un, ja tā, mēs atgriežam tās vērtÄ«bu. Mēs arÄ« pārbaudām, vai atslēga ir tukÅ”a, un, ja tā, mēs pārtraucam meklÄ“Å”anu.

Ja nevaram atrast atslēgu, kods atgriež tukÅ”u vērtÄ«bu.

Visas Ŕīs meklÄ“Å”anas darbÄ«bas var veikt vienlaikus, ievietojot un dzÄ“Å”ot. Katram tabulas pārim bÅ«s viens no četriem iepriekÅ” aprakstÄ«tajiem plÅ«smas stāvokļiem.

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 tiek veikta neparastā veidā: mēs atstājam atslēgu tabulā un atzÄ«mējam tās vērtÄ«bu (nevis paÅ”u atslēgu) kā tukÅ”u. Å is kods ir ļoti lÄ«dzÄ«gs lookup(), izņemot to, ka, ja atslēgai tiek atrasta atbilstÄ«ba, tās vērtÄ«ba tiek tukÅ”a.

Kā minēts iepriekÅ”, kad atslēga ir ierakstÄ«ta slotā, tā vairs netiek pārvietota. Pat tad, kad elements tiek izdzēsts no tabulas, atslēga paliek vietā, tā vērtÄ«ba vienkārÅ”i kļūst tukÅ”a. Tas nozÄ«mē, ka slota vērtÄ«bai nav jāizmanto atomu rakstÄ«Å”anas operācija, jo nav nozÄ«mes, vai paÅ”reizējā vērtÄ«ba ir tukÅ”a vai nē ā€“ tā tik un tā kļūs tukÅ”a.

Jauktabulas izmēra maiņa

Hash tabulas izmēru var mainÄ«t, izveidojot lielāku tabulu un ievietojot tajā netukÅ”us elementus no vecās tabulas. Es neieviesu Å”o funkcionalitāti, jo vēlējos, lai koda paraugs bÅ«tu vienkārÅ”s. Turklāt CUDA programmās atmiņas pieŔķirÅ”ana bieži tiek veikta resursdatora kodā, nevis CUDA kodolā.

Rakstā JaukÅ”anas galds bez bloÄ·Ä“Å”anas un gaidÄ«Å”anas apraksta, kā modificēt Ŕādu ar bloÄ·Ä“Å”anu aizsargātu datu struktÅ«ru.

Konkurētspēja

IepriekÅ” minētajos funkciju koda fragmentos gpu_hashtable_insert(), _lookup() Šø _delete() vienlaikus apstrādāt vienu atslēgu-vērtÄ«bu pāri. Un zemāk gpu_hashtable_insert(), _lookup() Šø _delete() paralēli apstrādāt pāru masÄ«vu, katru pāri atseviŔķā GPU izpildes pavedienā:

// 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);
    }
}

BloÄ·Ä“Å”anas izturÄ«ga jaucēj tabula atbalsta vienlaicÄ«gus ievietoÅ”anu, uzmeklÄ“Å”anu un dzÄ“Å”anu. Tā kā atslēgu-vērtÄ«bu pāri vienmēr atrodas vienā no četriem stāvokļiem un atslēgas nepārvietojas, tabula garantē pareizÄ«bu pat tad, ja vienlaikus tiek izmantotas dažāda veida darbÄ«bas.

Tomēr, ja mēs paralēli apstrādāsim ievietoÅ”anas un dzÄ“Å”anas partiju un ja ievades pāru masÄ«vā ir dublētās atslēgas, mēs nevarēsim paredzēt, kuri pāri ā€œuzvarēsā€ ā€” tie tiks ierakstÄ«ti jaukÅ”anas tabulā pēdējie. Pieņemsim, ka mēs nosaucām ievietoÅ”anas kodu ar ievades pāru masÄ«vu A/0 B/1 A/2 C/3 A/4. Kad kods ir pabeigts, tiek izveidoti pāri B/1 Šø C/3 ir garantēti klāt tabulā, bet tajā paŔā laikā jebkurÅ” no pāriem parādÄ«sies tajā A/0, A/2 vai A/4. Tas var bÅ«t vai nevar bÅ«t problēma ā€” tas viss ir atkarÄ«gs no lietojumprogrammas. Iespējams, jau iepriekÅ” zināt, ka ievades masÄ«vā nav atslēgu dublikātu, vai arÄ« jums var bÅ«t vienalga, kura vērtÄ«ba tika ierakstÄ«ta pēdējā.

Ja jums tā ir problēma, jums ir jāsadala dublikātu pāri dažādos CUDA sistēmas izsaukumos. Programmā CUDA jebkura darbÄ«ba, kas izsauc kodolu, vienmēr tiek pabeigta pirms nākamā kodola izsaukuma (vismaz viena pavediena ietvaros. Dažādos pavedienos kodoli tiek izpildÄ«ti paralēli). IepriekÅ” minētajā piemērā, ja izsaucat vienu kodolu ar A/0 B/1 A/2 C/3, bet otrs ar A/4, tad atslēga A iegÅ«s vērtÄ«bu 4.

Tagad parunāsim par to, vai funkcijām vajadzētu bÅ«t lookup() Šø delete() izmantojiet vienkārÅ”u vai nepastāvÄ«gu rādÄ«tāju uz pāru masÄ«vu hash tabulā. CUDA dokumentācija Nosaka, ka:

Kompilators var izvēlēties optimizēt lasÄ«Å”anu un rakstÄ«Å”anu globālajā vai koplietotajā atmiņā... Å Ä«s optimizācijas var atspējot, izmantojot atslēgvārdu volatile: ... jebkura atsauce uz Å”o mainÄ«go tiek apkopota reālā atmiņas lasÄ«Å”anas vai rakstÄ«Å”anas instrukcijā.

PareizÄ«bas apsvērumi nav jāpiemēro volatile. Ja izpildes pavediens izmanto keÅ”atmiņā saglabātu vērtÄ«bu no agrākas lasÄ«Å”anas darbÄ«bas, tad tajā tiks izmantota nedaudz novecojusi informācija. Bet tomēr Ŕī ir informācija no pareizā hash tabulas stāvokļa noteiktā kodola izsaukuma brÄ«dÄ«. Ja jums ir jāizmanto jaunākā informācija, varat izmantot indeksu volatile, bet tad veiktspēja nedaudz samazināsies: pēc maniem testiem, dzÄ“Å”ot 32 miljonus elementu, ātrums samazinājās no 500 miljoniem dzÄ“Å”anas/sek lÄ«dz 450 miljoniem dzÄ“Å”anas/sek.

ŠŸŃ€Š¾ŠøŠ·Š²Š¾Š“ŠøтŠµŠ»ŃŒŠ½Š¾ŃŃ‚ŃŒ

Pārbaudē par 64 miljonu elementu ievietoÅ”anu un 32 miljonu no tiem dzÄ“Å”anu notiek konkurence starp std::unordered_map un GPU praktiski nav hash tabulas:

VienkārÅ”a jaucēj tabula GPU
std::unordered_map pavadÄ«ja 70 691 ms, ievietojot un izņemot elementus un pēc tam tos atbrÄ«vojot unordered_map (atbrÄ«voÅ”anās no miljoniem elementu aizņem daudz laika, jo iekŔā unordered_map tiek veiktas vairākas atmiņas pieŔķirÅ”anas). GodÄ«gi sakot, std:unordered_map pilnÄ«gi atŔķirÄ«gi ierobežojumi. Tas ir viens CPU izpildes pavediens, atbalsta jebkura lieluma atslēgas vērtÄ«bas, labi darbojas ar augstu izmantoÅ”anas lÄ«meni un uzrāda stabilu veiktspēju pēc vairākkārtējas dzÄ“Å”anas.

Hash tabulas ilgums GPU un starpprogrammu komunikācijai bija 984 ms. Tas ietver laiku, kas pavadÄ«ts, ievietojot tabulu atmiņā un dzÄ“Å”ot to (vienreiz tiek pieŔķirts 1 GB atmiņas, kas CUDA aizņem kādu laiku), ievietojot un dzÄ“Å”ot elementus, kā arÄ« atkārtojot tos. Tiek ņemtas vērā arÄ« visas kopijas uz videokartes atmiņu un no tās.

PaÅ”as jaucēj tabulas aizpildÄ«Å”anai bija nepiecieÅ”ams 271 ms. Tas ietver laiku, ko videokarte pavada elementu ievietoÅ”anai un dzÄ“Å”anai, un neņem vērā laiku, kas pavadÄ«ts kopÄ“Å”anai atmiņā un atkārtoÅ”anai iegÅ«tajā tabulā. Ja GPU tabula darbojas ilgu laiku vai ja jaukÅ”anas tabula pilnÄ«bā atrodas videokartes atmiņā (piemēram, lai izveidotu jaucējtabulu, kuru izmantos cits GPU kods, nevis centrālais procesors), tad testa rezultāts ir bÅ«tisks.

Videokartes hash tabula demonstrē augstu veiktspēju, pateicoties lielai caurlaidspējai un aktīvai paralēlizācijai.

Ierobežojumi

Jaucēj tabulas arhitektūrā ir dažas problēmas, kas jāņem vērā:

  • Lineāro zondÄ“Å”anu apgrÅ«tina klasterizācija, kuras dēļ taustiņi tabulā nav novietoti nevainojami.
  • Izmantojot Å”o funkciju, taustiņi netiek noņemti delete un laika gaitā tie pārblÄ«vē galdu.

Tā rezultātā jaucēj tabulas veiktspēja var pakāpeniski pasliktināties, it Ä«paÅ”i, ja tā pastāv ilgu laiku un tajā ir daudz ievietoÅ”anas un dzÄ“Å”anas. Viens no veidiem, kā mazināt Å”os trÅ«kumus, ir pārjaukÅ”ana jaunā tabulā ar diezgan zemu izmantoÅ”anas lÄ«meni un atkārtotas jaukÅ”anas laikā filtrēt noņemtās atslēgas.

Lai ilustrētu aprakstÄ«tās problēmas, es izmantoÅ”u iepriekÅ” minēto kodu, lai izveidotu tabulu ar 128 miljoniem elementu un cilpu cauri 4 miljoniem elementu, lÄ«dz bÅ«Å”u aizpildÄ«jis 124 miljonus laika niÅ”u (izmantoÅ”anas lÄ«menis aptuveni 0,96). Å eit ir rezultātu tabula, katra rinda ir CUDA kodola izsaukums, lai vienā hash tabulā ievietotu 4 miljonus jaunu elementu:

LietoŔanas līmenis
IevietoŔanas ilgums 4 194 304 elementi

0,00
11,608448 ms (361,314798 miljoni atslēgu/sek.)

0,03
11,751424 ms (356,918799 miljoni atslēgu/sek.)

0,06
11,942592 ms (351,205515 miljoni atslēgu/sek.)

0,09
12,081120 ms (347,178429 miljoni atslēgu/sek.)

0,12
12,242560 ms (342,600233 miljoni atslēgu/sek.)

0,16
12,396448 ms (338,347235 miljoni atslēgu/sek.)

0,19
12,533024 ms (334,660176 miljoni atslēgu/sek.)

0,22
12,703328 ms (330,173626 miljoni atslēgu/sek.)

0,25
12,884512 ms (325,530693 miljoni atslēgu/sek.)

0,28
13,033472 ms (321,810182 miljoni atslēgu/sek.)

0,31
13,239296 ms (316,807174 miljoni atslēgu/sek.)

0,34
13,392448 ms (313,184256 miljoni atslēgu/sek.)

0,37
13,624000 ms (307,861434 miljoni atslēgu/sek.)

0,41
13,875520 ms (302,280855 miljoni atslēgu/sek.)

0,44
14,126528 ms (296,909756 miljoni atslēgu/sek.)

0,47
14,399328 ms (291,284699 miljoni atslēgu/sek.)

0,50
14,690304 ms (285,515123 miljoni atslēgu/sek.)

0,53
15,039136 ms (278,892623 miljoni atslēgu/sek.)

0,56
15,478656 ms (270,973402 miljoni atslēgu/sek.)

0,59
15,985664 ms (262,379092 miljoni atslēgu/sek.)

0,62
16,668673 ms (251,627968 miljoni atslēgu/sek.)

0,66
17,587200 ms (238,486174 miljoni atslēgu/sek.)

0,69
18,690048 ms (224,413765 miljoni atslēgu/sek.)

0,72
20,278816 ms (206,831789 miljoni atslēgu/sek.)

0,75
22,545408 ms (186,038058 miljoni atslēgu/sek.)

0,78
26,053312 ms (160,989275 miljoni atslēgu/sek.)

0,81
31,895008 ms (131,503463 miljoni atslēgu/sek.)

0,84
42,103294 ms (99,619378 miljoni atslēgu/sek.)

0,87
61,849056 ms (67,815164 miljoni atslēgu/sek.)

0,90
105,695999 ms (39,682713 miljoni atslēgu/sek.)

0,94
240,204636 ms (17,461378 miljoni atslēgu/sek.)

Palielinoties izmantoÅ”anai, veiktspēja samazinās. Vairumā gadÄ«jumu tas nav vēlams. Ja lietojumprogramma ievieto elementus tabulā un pēc tam tos izmet (piemēram, saskaitot vārdus grāmatā), tad tā nav problēma. Bet, ja lietojumprogramma izmanto ilgstoÅ”as ā€‹ā€‹ā€‹ā€‹jaucējtabulas (piemēram, grafiskajā redaktorā, lai saglabātu netukÅ”as attēlu daļas, kurās lietotājs bieži ievieto un dzÄ“Å” informāciju), Ŕī darbÄ«ba var bÅ«t problemātiska.

Un izmērÄ«ja hash tabulas zondÄ“Å”anas dziļumu pēc 64 miljoniem ievietoÅ”anas (izmantoÅ”anas koeficients 0,5). Vidējais dziļums bija 0,4774, tāpēc lielākā daļa taustiņu atradās vai nu labākajā iespējamajā slotā, vai viena slota attālumā no labākās pozÄ«cijas. Maksimālais zondÄ“Å”anas dziļums bija 60.

Pēc tam es mērÄ«ju zondÄ“Å”anas dziļumu uz galda ar 124 miljoniem ieliktņu (izmantoÅ”anas koeficients 0,97). Vidējais dziļums jau bija 10,1757, bet maksimālais - 6474 (!!). Lineārās sensora veiktspēja ievērojami samazinās pie augsta izmantoÅ”anas lÄ«meņa.

Vislabāk ir saglabāt zemu Ŕīs hash tabulas izmantoÅ”anas lÄ«meni. Bet tad mēs palielinām veiktspēju uz atmiņas patēriņa rēķina. Par laimi, 32 bitu atslēgu un vērtÄ«bu gadÄ«jumā to var attaisnot. Ja iepriekÅ” minētajā piemērā tabulā ar 128 miljoniem elementu saglabājam izmantoÅ”anas koeficientu 0,25, tad tajā varam ievietot ne vairāk kā 32 miljonus elementu, un tiks zaudēti atlikuÅ”ie 96 miljoni slotu - 8 baiti katram pārim. , 768 MB zaudētas atmiņas.

LÅ«dzu, ņemiet vērā, ka mēs runājam par videokartes atmiņas zudumu, kas ir vērtÄ«gāks resurss nekā sistēmas atmiņa. Lai gan lielākajai daļai mÅ«sdienu galddatoru grafisko karÅ”u, kas atbalsta CUDA, ir vismaz 4 GB atmiņa (raksta tapÅ”anas brÄ«dÄ« NVIDIA 2080 Ti ir 11 GB), zaudēt Ŕādas summas tomēr nebÅ«tu prātÄ«gākais lēmums.

Vēlāk rakstÄ«Å”u vairāk par jaucējtabulu izveidi videokartēm, kurām nav problēmu ar zondÄ“Å”anas dziļumu, kā arÄ« veidiem, kā atkārtoti izmantot izdzēstos slotus.

Zondes dziļuma mērÄ«Å”ana

Lai noteiktu atslēgas pārbaudes dziļumu, mēs varam iegūt atslēgas jaucējkodu (tā ideālo tabulas indeksu) no tās faktiskā tabulas indeksa:

// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);

Tā kā divi divi komplementa binārie skaitļi ir burvÄ«gi un ka jaukÅ”anas tabulas ietilpÄ«ba ir divi ar pakāpju divi, Ŕī pieeja darbosies pat tad, ja atslēgas indekss tiek pārvietots uz tabulas sākumu. Paņemsim atslēgu, kas ir sajaukta ar 1, bet tiek ievietota slotā 3. Tad tabulai ar ietilpÄ«bu 4 mēs iegÅ«stam (3 ā€” 1) & 3, kas ir lÄ«dzvērtÄ«gs 2.

Secinājums

Ja jums ir jautājumi vai komentāri, lūdzu, rakstiet man uz e-pastu Twitter vai atveriet jaunu tēmu krātuves.

Šis kods tika uzrakstīts, iedvesmojoties no lieliskiem rakstiem:

Nākotnē es turpināŔu rakstÄ«t par video karÅ”u hash tabulu ievieÅ”anu un analizēt to veiktspēju. Manos plānos ietilpst ķēde, Robina Huda jaukÅ”ana un dzeguzes jaukÅ”ana, izmantojot atomu operācijas datu struktÅ«rās, kas ir draudzÄ«gas GPU.

Avots: www.habr.com

Pievieno komentāru