Ndidalemba pa Github pulojekiti yatsopano A Simple GPU Hash Table.

Ndi tebulo losavuta la GPU hashi lomwe limatha kukonza mamiliyoni mazana oyika pamphindikati. Pa laputopu yanga ya NVIDIA GTX 1060, code imayika ma 64 miliyoni omwe amapangidwa mwachisawawa pafupifupi 210 ms ndikuchotsa ma 32 miliyoni awiriawiri pafupifupi 64 ms.

Ndiko kuti, liwiro pa laputopu ndi pafupifupi 300 miliyoni amaika/sekondi ndi 500 miliyoni deletes/mphindi.

Gome lalembedwa mu CUDA, ngakhale njira yomweyo ingagwiritsidwe ntchito ku HLSL kapena GLSL. Kukhazikitsa kuli ndi malire angapo kuti muwonetsetse kuti pakompyuta ikugwira ntchito kwambiri:

Makiyi a 32-bit okha ndi zofananira zomwe zimasinthidwa.
Tebulo la hashi lili ndi kukula kokhazikika.
Ndipo kukula uku kuyenera kukhala kofanana ndi ziwiri ku mphamvu.

Pa makiyi ndi zofunikira, muyenera kusungitsa cholembera chosavuta (mu code yomwe ili pamwambapa iyi ndi 0xffffffff).

Tebulo la hashi lopanda maloko

Tebulo la hashi limagwiritsa ntchito ma adilesi otseguka ndi kufufuza kwa mzere, ndiye kuti, ndi mndandanda wamagulu awiri amtengo wapatali omwe amasungidwa pamtima ndipo ali ndi ntchito yabwino kwambiri ya cache. Zomwezo sizinganenedwe pa unyolo, womwe umaphatikizapo kufunafuna cholozera pamndandanda wolumikizidwa. Gome la hashi ndi njira yosavuta yosungira zinthu KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Kukula kwa tebulo ndi mphamvu ziwiri, osati nambala yaikulu, chifukwa malangizo amodzi ofulumira ndi okwanira kugwiritsa ntchito pow2 / AND mask, koma woyendetsa modulus amachedwa kwambiri. Izi ndizofunikira pakufufuza kwa mzere, chifukwa poyang'ana patebulo lokhala ndi mzere wolozera uyenera kukulungidwa pagawo lililonse. Ndipo chifukwa chake, mtengo wa ntchitoyo umawonjezedwa modulo pagawo lililonse.

Tebulo limangosunga makiyi ndi mtengo wa chinthu chilichonse, osati makiyi. Popeza tebulo limangosunga makiyi a 32-bit, hashi imawerengedwa mwachangu kwambiri. Khodi yomwe ili pamwambapa imagwiritsa ntchito hashi ya Murmur3, yomwe imangosintha pang'ono, ma XOR ndi kuchulukitsa.

Tebulo la hashi limagwiritsa ntchito njira zotetezera zotsekera zomwe sizimayenderana ndi dongosolo la kukumbukira. Ngakhale zolemba zina zitasokoneza dongosolo la machitidwe ena otere, tebulo la hashi likhalabe lolondola. Tikambirana za izi pansipa. Njirayi imagwira ntchito bwino ndi makadi amakanema omwe amayendetsa ulusi masauzande nthawi imodzi.

Makiyi ndi zofunikira pa tebulo la hashi zimayambitsidwa kuti zisakhale zopanda kanthu.

Khodiyo imatha kusinthidwa kuti igwire makiyi a 64-bit komanso mfundo zake. Makiyi amafunikira ma atomiki owerengera, kulemba, ndi kufananiza-ndi-kusinthana ntchito. Ndipo zofunikira zimafunikira ma atomiki owerengera ndi kulemba. Mwamwayi, mu CUDA, ntchito zowerengera-zolemba za 32- ndi 64-bit values ndi atomiki bola zimagwirizana mwachilengedwe (onani pansipa). apa), ndi makhadi amakono amakono amathandizira 64-bit atomic kuyerekezera-ndi-kusinthanitsa ntchito. Zachidziwikire, mukasunthira ku 64 bits, magwiridwe antchito amachepera pang'ono.

Hash table state

Gulu lililonse lamtengo wapatali pa tebulo la hashi likhoza kukhala ndi chimodzi mwa zigawo zinayi:

Mfungulo ndi mtengo zilibe kanthu. Munthawi imeneyi, tebulo la hashi limayambitsidwa.
Chinsinsi chalembedwa, koma mtengo wake sunalembedwe. Ngati ulusi wina ukuwerenga deta pano, umabwerera wopanda kanthu. Izi ndizabwinobwino, zomwezo zikadachitika ngati ulusi wina wophatikizira udagwira ntchito pang'ono, ndipo tikukamba za kapangidwe ka data kamodzi.
Zonse fungulo ndi mtengo zimalembedwa.
Mtengowu ulipo ku ulusi wina wophatikizika, koma chinsinsi sichinafike. Izi zitha kuchitika chifukwa mtundu wamapulogalamu a CUDA uli ndi mtundu wa kukumbukira wosasamala. Izi ndizabwinobwino; mulimonse, fungulo likadalibe, ngakhale mtengo suli choncho.

Chofunikira ndichakuti fungulo likangolembedwa ku slot, silisunthanso - ngakhale fungulo litachotsedwa, tikambirana pansipa.

Khodi ya tebulo la hashi imagwiranso ntchito ndi mitundu yokumbukira mosasamala momwe momwe kukumbukira kumawerengedwera ndi kulembedwa sikudziwika. Pamene tikuyang'ana kuyika, kuyang'ana, ndi kuchotsa mu tebulo la hashi, kumbukirani kuti gulu lirilonse lamtengo wapatali liri m'modzi mwa zigawo zinayi zomwe zafotokozedwa pamwambapa.

Kuyika mu tebulo la hashi

Ntchito ya CUDA yomwe imayika awiriawiri amtengo wapatali pa tebulo la hashi imawoneka motere:

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

Kuti muyike kiyi, kachidindo kameneka kamadutsa mumndandanda wa tebulo la hashi kuyambira ndi hashi ya kiyi yoyikidwa. Kagawo kalikonse kagawo kameneka kamagwira ntchito yofananitsa ndi kusinthana kwa atomiki yomwe imafanizira makiyi omwe ali mugawolo ndi opanda kanthu. Ngati zindikirani kuti palibe kugwirizana, kiyi mu slot imasinthidwa ndi kiyi yoyikidwa, ndiyeno kiyi yoyambira imabwezeretsedwa. Ngati kiyi yoyambirirayi inali yopanda kanthu kapena ikufanana ndi kiyi yomwe idayikidwapo, ndiye kuti codeyo idapeza malo oyenera kuyikapo ndikuyika mtengo womwe adalowetsedwamo.

Ngati muyimba kernel imodzi gpu_hashtable_insert() pali zinthu zingapo zokhala ndi fungulo lomwelo, ndiye kuti zilizonse zomwe zimafunikira zitha kulembedwa pakiyi. Izi zimawonedwa ngati zachilendo: chimodzi mwazofunikira zomwe zimalemba pakuyimba foni zitha kuchita bwino, koma popeza zonsezi zimachitika limodzi mkati mwa mizere ingapo yophatikizika, sitingathe kulosera kuti ndi kukumbukira kotani komwe kudzakhala komaliza.

Kufufuza kwa tebulo la hash

Khodi yosaka makiyi:

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

Kuti tipeze mtengo wa kiyi yosungidwa patebulo, timabwereza mndandanda kuyambira ndi hashi ya kiyi yomwe tikuyang'ana. Pagawo lililonse, timayang'ana ngati fungulo ndilomwe tikuyang'ana, ndipo ngati ndi choncho, timabwezera mtengo wake. Timayang'ananso ngati kiyiyo ilibe kanthu, ndipo ngati ndi choncho, timachotsa kusaka.

Ngati sitingapeze fungulo, code imabwezeretsa mtengo wopanda kanthu.

Zosaka zonsezi zitha kuchitidwa nthawi imodzi kudzera muzoyikapo ndi kuchotsa. Gulu lirilonse patebulo lidzakhala ndi limodzi mwa zigawo zinayi zomwe zafotokozedwa pamwambapa za kuyenda.

Kuchotsa mu tebulo la hashi

Khodi yochotsa makiyi:

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

Kuchotsa fungulo kumachitika mwachilendo: timasiya fungulo patebulo ndikulemba mtengo wake (osati fungulo lokha) lopanda kanthu. Code iyi ndi yofanana kwambiri ndi lookup(), kupatula kuti machesi akapezeka pa kiyi, amapangitsa kuti mtengo wake ukhale wopanda kanthu.

Monga tafotokozera pamwambapa, kiyi ikangolembedwa ku slot, sisunthanso. Ngakhale chinthu chikachotsedwa patebulo, fungulo limakhalabe m'malo, mtengo wake umakhala wopanda kanthu. Izi zikutanthauza kuti sitiyenera kugwiritsa ntchito ma atomiki kulemba pa mtengo wa slot, chifukwa zilibe kanthu kaya mtengo wapano uli wopanda kanthu kapena ayi - ukhalabe wopanda kanthu.

Kusintha kukula kwa tebulo la hashi

Mutha kusintha kukula kwa tebulo la hashi popanga tebulo lalikulu ndikuyika zinthu zopanda kanthu kuchokera patebulo lakale momwemo. Sindinagwiritse ntchito izi chifukwa ndinkafuna kuti chitsanzochi chikhale chosavuta. Komanso, m'mapulogalamu a CUDA, kugawa kukumbukira nthawi zambiri kumachitika mu code host osati mu CUDA kernel.

M'nkhaniyi Table ya Hashi Yopanda Lokhoma limafotokoza momwe mungasinthire deta yotetezedwa ndi loko.

Kupikisana

Pazidutswa tambiri zamakhodi a ntchito gpu_hashtable_insert(), _lookup() и _delete() sinthani mtengo umodzi wamtengo umodzi nthawi imodzi. Ndipo pansi gpu_hashtable_insert(), _lookup() и _delete() sungani magulu awiriawiri mofanana, gulu lililonse mu ulusi wotsatira wa GPU:

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

Gome la hashi losatsekeka limathandizira kuyikapo nthawi imodzi, kuyang'ana, ndi kuchotsa. Chifukwa mawiri amtengo wapatali nthawi zonse amakhala m'modzi mwa zigawo zinayi ndipo makiyi sasuntha, tebulo limatsimikizira kulondola ngakhale mitundu yosiyanasiyana ya ntchito ikugwiritsidwa ntchito nthawi imodzi.

Komabe, ngati tikonza zoyikapo ndi zochotsa mofananira, ndipo ngati zolowetsazo zili ndi makiyi obwereza, ndiye kuti sitingathe kulosera kuti ndi mapeyala ati "adzapambana" -adzalembedwa ku tebulo lomaliza. Tiyerekeze kuti tayitanitsa nambala yoyika ndi mitundu yolowera ya awiriawiri A/0 B/1 A/2 C/3 A/4. Code ikamaliza, awiriawiri B/1 и C/3 akutsimikiziridwa kukhalapo mu tebulo, koma nthawi yomweyo aliyense wa awiriawiri adzaoneka mmenemo A/0, A/2 kapena A/4. Izi zitha kukhala zovuta kapena sizingakhale zovuta - zonse zimatengera kugwiritsa ntchito. Mutha kudziwiratu kuti palibe makiyi obwereza muzolowera, kapena mwina simusamala kuti ndi phindu liti lomwe linalembedwa pomaliza.

Ngati ili ndi vuto kwa inu, ndiye kuti muyenera kulekanitsa awiriwa awiriwa mumayendedwe osiyanasiyana a CUDA. Mu CUDA, ntchito iliyonse yomwe imayitana kernel imamaliza kuyimbanso kernel (osachepera mkati mwa ulusi umodzi. Mu ulusi wosiyana, maso amapangidwa mofanana). Mu chitsanzo pamwambapa, ngati muyimbira kernel imodzi ndi A/0 B/1 A/2 C/3, ndi winayo ndi A/4, kenako kiyi A adzapeza mtengo 4.

Tsopano tiyeni tikambirane ngati ntchito ayenera lookup() и delete() gwiritsani ntchito cholozera chowoneka bwino kapena chosasunthika kumagulu angapo patebulo la hashi. Zolemba za CUDA Ananena kuti:

Wopangayo atha kusankha kukhathamiritsa zowerenga ndikulemba kuti zikumbukiro zapadziko lonse lapansi kapena zogawana... Kukhathamiritsa uku kumatha kuzimitsidwa pogwiritsa ntchito mawu osakira. volatile...

Malingaliro olondola safuna kugwiritsa ntchito volatile. Ngati ulusi wopha umagwiritsa ntchito mtengo wosungidwa kuchokera kuzomwe zidawerengedwa kale, ndiye kuti ukugwiritsa ntchito chidziwitso chachikale. Koma komabe, ichi ndi chidziwitso chochokera kumalo olondola a tebulo la hashi panthawi inayake ya kuyimba kwa kernel. Ngati mukufuna kugwiritsa ntchito chidziwitso chatsopano, mutha kugwiritsa ntchito index volatile, koma ndiye ntchitoyo idzachepa pang'ono: molingana ndi mayesero anga, pochotsa zinthu 32 miliyoni, liwiro linatsika kuchokera ku 500 miliyoni kuchotsa / mphindi kufika pa 450 miliyoni kuchotsa / mphindi.

Kukonzekera

Pakuyesa kuyika zinthu 64 miliyoni ndikuchotsa 32 miliyoni, mpikisano pakati std::unordered_map ndipo palibe tebulo la hashi la GPU:

std::unordered_map adawononga 70 ms ndikuyika ndikuchotsa zinthu ndikuzimasula unordered_map (kuchotsa mamiliyoni azinthu kumatenga nthawi yambiri, chifukwa mkati unordered_map kugawa kwamakumbukidwe angapo kumapangidwa). Kunena zoona, std:unordered_map zoletsa zosiyana kwathunthu. Ndi ulusi umodzi wa CPU wophatikizika, umathandizira makiyi amtundu uliwonse, umagwira ntchito bwino pamagwiritsidwe apamwamba, ndikuwonetsa magwiridwe antchito pambuyo pochotsa kangapo.

Kutalika kwa tebulo la hashi la GPU ndi kulumikizana pakati pa mapulogalamu kunali 984 ms. Izi zikuphatikizapo nthawi yoyika tebulo mu kukumbukira ndikulichotsa (kugawa 1 GB ya kukumbukira nthawi imodzi, zomwe zimatenga nthawi mu CUDA), kuyika ndi kuchotsa zinthu, ndi kubwereza. Makope onse opita ndi kuchokera pamtima pamakhadi amakanema amaganiziridwanso.

Tebulo la hashi palokha lidatenga 271 ms kuti limalize. Izi zikuphatikizanso nthawi yomwe khadi la kanema likulowetsa ndikuchotsa zinthu, ndipo sizitengera nthawi yomwe idagwiritsidwa ntchito kukopera pamtima ndikubwerezanso patebulo lomwe likubwera. Ngati tebulo la GPU limakhala nthawi yayitali, kapena ngati tebulo la hashi liri mu kukumbukira khadi la kanema (mwachitsanzo, kupanga tebulo la hashi lomwe lidzagwiritsidwe ntchito ndi ma code ena a GPU osati purosesa yapakati), ndiye zotsatira za mayeso ndi zogwirizana.

Tebulo la hashi la khadi la kanema likuwonetsa magwiridwe antchito apamwamba chifukwa cha kuchuluka kwamphamvu komanso kufanana kogwira ntchito.

zolakwa

Mapangidwe a tebulo la hash ali ndi zovuta zingapo zomwe muyenera kuzidziwa:

Kufufuza kwa mzere kumalepheretsedwa ndi kusanjana, zomwe zimapangitsa makiyi omwe ali patebulo kuyikidwa mocheperapo kuposa mwangwiro.
Makiyi samachotsedwa pogwiritsa ntchito ntchitoyi delete ndipo m’kupita kwa nthawi amasokoneza tebulo.

Zotsatira zake, magwiridwe antchito a tebulo la hashi amatha kutsika pang'onopang'ono, makamaka ngati atakhalapo kwa nthawi yayitali ndipo ali ndi zoyika zambiri ndikuchotsa. Njira imodzi yochepetsera kuipa kumeneku ndikubwezeretsanso patebulo latsopano ndikugwiritsa ntchito pang'onopang'ono ndikusefa makiyi omwe adachotsedwa pakukonzanso.

Kuti ndiwonetsere zomwe tafotokozazi, ndigwiritsa ntchito nambala yomwe ili pamwambapa kuti ndipange tebulo lokhala ndi zinthu 128 miliyoni ndikudumphadumpha 4 miliyoni mpaka nditadzaza mipata 124 miliyoni (magwiritsidwe ntchito pafupifupi 0,96). Nayi tebulo lotsatira, mzere uliwonse ndi kuyimba kwa CUDA kuti muyike zinthu zatsopano 4 miliyoni patebulo limodzi la hashi:

Mtengo wogwiritsa ntchito
Kutalika kwazinthu 4

0,00
11,608448 ms (361,314798 miliyoni makiyi/mphindi.)

0,03
11,751424 ms (356,918799 miliyoni makiyi/mphindi.)

0,06
11,942592 ms (351,205515 miliyoni makiyi/mphindi.)

0,09
12,081120 ms (347,178429 miliyoni makiyi/mphindi.)

0,12
12,242560 ms (342,600233 miliyoni makiyi/mphindi.)

0,16
12,396448 ms (338,347235 miliyoni makiyi/mphindi.)

0,19
12,533024 ms (334,660176 miliyoni makiyi/mphindi.)

0,22
12,703328 ms (330,173626 miliyoni makiyi/mphindi.)

0,25
12,884512 ms (325,530693 miliyoni makiyi/mphindi.)

0,28
13,033472 ms (321,810182 miliyoni makiyi/mphindi.)

0,31
13,239296 ms (316,807174 miliyoni makiyi/mphindi.)

0,34
13,392448 ms (313,184256 miliyoni makiyi/mphindi.)

0,37
13,624000 ms (307,861434 miliyoni makiyi/mphindi.)

0,41
13,875520 ms (302,280855 miliyoni makiyi/mphindi.)

0,44
14,126528 ms (296,909756 miliyoni makiyi/mphindi.)

0,47
14,399328 ms (291,284699 miliyoni makiyi/mphindi.)

0,50
14,690304 ms (285,515123 miliyoni makiyi/mphindi.)

0,53
15,039136 ms (278,892623 miliyoni makiyi/mphindi.)

0,56
15,478656 ms (270,973402 miliyoni makiyi/mphindi.)

0,59
15,985664 ms (262,379092 miliyoni makiyi/mphindi.)

0,62
16,668673 ms (251,627968 miliyoni makiyi/mphindi.)

0,66
17,587200 ms (238,486174 miliyoni makiyi/mphindi.)

0,69
18,690048 ms (224,413765 miliyoni makiyi/mphindi.)

0,72
20,278816 ms (206,831789 miliyoni makiyi/mphindi.)

0,75
22,545408 ms (186,038058 miliyoni makiyi/mphindi.)

0,78
26,053312 ms (160,989275 miliyoni makiyi/mphindi.)

0,81
31,895008 ms (131,503463 miliyoni makiyi/mphindi.)

0,84
42,103294 ms (99,619378 miliyoni makiyi/mphindi.)

0,87
61,849056 ms (67,815164 miliyoni makiyi/mphindi.)

0,90
105,695999 ms (39,682713 miliyoni makiyi/mphindi.)

0,94
240,204636 ms (17,461378 miliyoni makiyi/mphindi.)

Pamene kugwiritsidwa ntchito kukuwonjezeka, ntchito imachepa. Izi sizofunika nthawi zambiri. Ngati pulogalamuyo iyika zinthu patebulo ndikuzitaya (mwachitsanzo, powerenga mawu m'buku), ndiye kuti ili si vuto. Koma ngati pulogalamuyo imagwiritsa ntchito tebulo la hashi lalitali (mwachitsanzo, muzojambula zojambula kuti musunge zithunzi zopanda kanthu zomwe wogwiritsa ntchito nthawi zambiri amaika ndikuchotsa zambiri), ndiye kuti khalidweli likhoza kukhala lovuta.

Ndipo anayeza tebulo la hashi loyesa kuya pambuyo pakuyika 64 miliyoni (magwiritsidwe 0,5). Kuzama kwapakati kunali 0,4774, kotero makiyi ambiri anali m'malo abwino kwambiri kapena malo amodzi kutali ndi malo abwino kwambiri. Kuzama kwakukulu kwa mawu kunali 60.

Kenako ndidayezera kuya kwake patebulo yokhala ndi zoyika 124 miliyoni (magwiritsidwe 0,97). Kuzama kwapakati kunali kale 10,1757, ndipo pazipita - 6474 (!!). Kuchita kwa ma Linear sensing kumatsika kwambiri pakugwiritsa ntchito kwambiri.

Ndibwino kuti muchepetse kugwiritsa ntchito tebulo la hashi. Koma ndiye timawonjezera magwiridwe antchito pakuwononga kukumbukira. Mwamwayi, pankhani ya makiyi a 32-bit ndi mfundo zake, izi zitha kulungamitsidwa. Ngati mu chitsanzo pamwambapa, patebulo lokhala ndi zinthu 128 miliyoni, timasunga zomwe zimagwiritsidwa ntchito ndi 0,25, ndiye kuti sitingathe kuyikamo zinthu zopitilira 32 miliyoni, ndipo mipata yotsalayo 96 miliyoni idzatayika - ma byte 8 pa gulu lililonse. , 768 MB ya kukumbukira kotayika.

Chonde dziwani kuti tikukamba za kutayika kwa kukumbukira khadi ya kanema, yomwe ndi chinthu chofunika kwambiri kuposa kukumbukira dongosolo. Ngakhale makhadi ambiri amakono apakompyuta omwe amathandizira CUDA ali ndi kukumbukira osachepera 4 GB (panthawi yolemba, NVIDIA 2080 Ti ili ndi 11 GB), sikungakhale chisankho chanzeru kutaya ndalama zotere.

Pambuyo pake ndilemba zambiri za kupanga matebulo a hashi a makadi amakanema omwe alibe vuto pakufufuza mwakuya, komanso njira zogwiritsiranso ntchito mipata yochotsedwa.

Muyeso wakuya womveka

Kuti tidziwe kuzama kwa kiyi, titha kuchotsa hashi ya kiyiyo (mlozera wabwino wa tebulo) kuchokera pa tebulo lake lenileni:

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

Chifukwa chamatsenga aawiri awiri omwe amathandizira manambala a binary komanso kuti mphamvu ya tebulo la hashi ndi yachiwiri ku mphamvu ya awiri, njira iyi idzagwira ntchito ngakhale pamene mndandanda wachinsinsi umasunthira kumayambiriro kwa tebulo. Tiyeni titenge kiyi yomwe imathamangira ku 1, koma imayikidwa mu slot 3. Kenako patebulo lokhala ndi mphamvu 4 timapeza. (3 — 1) & 3, yomwe ikufanana ndi 2.

Pomaliza

Ngati muli ndi mafunso kapena ndemanga, chonde nditumizireni imelo Twitter kapena tsegulani mutu watsopano nkhokwe.

Khodi iyi idalembedwa mowuziridwa ndi zolemba zabwino kwambiri:

M'tsogolomu, ndipitiriza kulemba za kukhazikitsidwa kwa tebulo la hashi pamakhadi a kanema ndikuwunika momwe amachitira. Zolinga zanga zikuphatikiza unyolo, Robin Hood hashing, ndi cuckoo hashing pogwiritsa ntchito ma atomiki pamapangidwe a data omwe ali ochezeka ndi GPU.

Source: www.habr.com

Gome losavuta la hashi la GPU