Clàr hash sìmplidh airson GPU

Clàr hash sìmplidh airson GPU
Chuir mi suas e air Github pròiseact ùr Clàr Hash GPU sìmplidh.

Is e clàr hash GPU sìmplidh a th’ ann a tha comasach air na ceudan de mhilleanan de chuir a-steach gach diog a ghiullachd. Air an laptop NVIDIA GTX 1060 agam, bidh an còd a’ cuir a-steach 64 millean paidhrichean prìomh luach air an cruthachadh air thuaiream ann an timcheall air 210 ms agus a’ toirt air falbh 32 millean paidhir ann an timcheall air 64 ms.

Is e sin, tha an astar air laptop timcheall air 300 millean cuir a-steach / diog agus 500 millean cuir às / diog.

Tha an clàr sgrìobhte ann an CUDA, ged a dh’ fhaodar an aon dòigh-obrach a chleachdadh airson HLSL no GLSL. Tha grunn chuingealachaidhean aig a’ bhuileachadh gus dèanamh cinnteach à àrd-choileanadh air cairt bhidio:

  • Chan eil ach iuchraichean 32-bit agus na h-aon luachan air an giullachd.
  • Tha meud stèidhichte aig a’ bhòrd hash.
  • Agus feumaidh am meud seo a bhith co-ionann ri dhà ris a 'chumhachd.

Airson iuchraichean agus luachan, feumaidh tu comharraiche delimiter sìmplidh a ghlèidheadh ​​​​(anns a’ chòd gu h-àrd is e seo 0xffffffff).

Clàr hash gun ghlasan

Bidh an clàr hash a’ cleachdadh seòladh fosgailte le sgrùdadh loidhneach, is e sin, is e dìreach sreath de chàraidean prìomh luach a tha air an stòradh mar chuimhneachan agus aig a bheil coileanadh tasgadan nas fheàrr. Chan urrainnear an aon rud a ràdh airson slabhraidhean, a tha a’ toirt a-steach lorg puing ann an liosta ceangailte. Tha clàr hash na raon sìmplidh airson eileamaidean stòraidh KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Tha meud a’ bhùird na chumhachd de dhà, chan e prìomh àireamh, oir tha aon stiùireadh luath gu leòr airson am masg pow2/AND a chuir an sàs, ach tha an gnìomhaiche modulus tòrr nas slaodaiche. Tha seo cudromach a thaobh sgrùdadh sreathach, oir ann an clàr sreathach feumaidh an clàr-amais slot a bhith air a phasgadh anns gach slot. Agus mar thoradh air an sin, tha cosgais na h-obrach air a chur ris modulo anns gach slot.

Chan eil anns a’ chlàr ach an iuchair agus an luach airson gach eileamaid a stòradh, chan e hash den iuchair. Leis nach eil am bòrd a’ stòradh ach iuchraichean 32-bit, tha an hash air a thomhas gu math luath. Bidh an còd gu h-àrd a’ cleachdadh hash Murmur3, nach bi a’ coileanadh ach beagan ghluasadan, XORs agus iomadachadh.

Bidh an clàr hash a’ cleachdadh dòighean dìon glasaidh a tha neo-eisimeileach bho òrdugh cuimhne. Fiù ma bhios cuid de dh’ obraichean sgrìobhaidh a’ cur dragh air òrdugh obrachaidhean eile mar sin, cumaidh am bòrd hash an staid cheart fhathast. Bruidhnidh sinn mu dheidhinn seo gu h-ìosal. Bidh an dòigh-obrach ag obair gu math le cairtean bhidio a bhios a’ ruith mìltean de snàithleanan aig an aon àm.

Tha na h-iuchraichean agus na luachan sa chlàr hash air an tòiseachadh gu falamh.

Faodar an còd atharrachadh gus iuchraichean agus luachan 64-bit a làimhseachadh cuideachd. Feumaidh iuchraichean obair leughaidh, sgrìobhaidh, agus coimeas-is-iomlaid atamach. Agus feumaidh luachan obrachaidhean leughaidh is sgrìobhaidh atamach. Gu fortanach, ann an CUDA, tha gnìomhachd leughaidh-sgrìobhaidh airson luachan 32- agus 64-bit atamach fhad ‘s a tha iad air an co-thaobhadh gu nàdarra (faic gu h-ìosal). an seo), agus tha cairtean bhidio ùr-nodha a’ toirt taic do ghnìomhachd coimeas-is-iomlaid atamach 64-bit. Gu dearbh, nuair a ghluaiseas tu gu 64 pìosan, bidh coileanadh a’ dol sìos beagan.

Stàit clàr hash

Faodaidh aon de cheithir stàitean a bhith aig gach paidhir prìomh luach ann an clàr hash:

  • Tha iuchair agus luach falamh. Anns an stàit seo, tha an clàr hash air a thòiseachadh.
  • Chaidh an iuchair a sgrìobhadh sìos, ach cha deach an luach a sgrìobhadh fhathast. Ma tha snàithlean eile a’ leughadh dàta an-dràsta, tillidh e falamh. Tha seo àbhaisteach, bhiodh an aon rud air tachairt nam biodh snàithlean cur gu bàs eile air obrachadh beagan na bu thràithe, agus tha sinn a’ bruidhinn air structar dàta co-aontach.
  • Tha an dà chuid an iuchair agus an luach air an clàradh.
  • Tha an luach ri fhaighinn le snàithleanan eile de choileanadh, ach chan eil an iuchair fhathast. Faodaidh seo tachairt leis gu bheil modal cuimhne air òrdachadh gu soilleir aig modal prògramadh CUDA. Tha seo àbhaisteach; co-dhiù, tha an iuchair fhathast falamh, eadhon ged nach eil an luach mar sin tuilleadh.

Is e nuance cudromach, aon uair ‘s gu bheil an iuchair air a sgrìobhadh chun t-slot, nach gluais e tuilleadh - eadhon ged a thèid an iuchair a dhubhadh às, bruidhnidh sinn mu dheidhinn seo gu h-ìosal.

Bidh an còd clàr hash eadhon ag obair le modalan cuimhne òrdaichte anns nach eil fios dè an òrdugh anns a bheil cuimhne air a leughadh agus air a sgrìobhadh. Mar a choimheadas sinn air cuir a-steach, lorg, agus cuir às ann an clàr hash, cuimhnich gu bheil gach paidhir luach-iuchrach ann an aon de na ceithir stàitean a chaidh a mhìneachadh gu h-àrd.

A 'cur a-steach do bhòrd hash

Tha an gnìomh CUDA a chuireas a-steach paidhrichean prìomh luach a-steach do bhòrd hash a’ coimhead mar seo:

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

Gus iuchair a chuir a-steach, bidh an còd ag ath-aithris tron ​​​​chlàr hash table a’ tòiseachadh le hash na h-iuchrach a chaidh a chuir a-steach. Bidh gach slot san raon a’ coileanadh gnìomhachd coimeas-is-iomlaid atamach a nì coimeas eadar an iuchair san t-slot sin gu falamh. Ma lorgar mì-chothromachadh, thèid an iuchair san t-slot ùrachadh leis an iuchair a chaidh a chuir a-steach, agus an uairsin thèid an iuchair slot tùsail a thilleadh. Ma bha an iuchair thùsail seo falamh no ma bha e co-ionnan ris an iuchair a chaidh a chuir a-steach, lorg an còd slot iomchaidh airson a chuir a-steach agus chuir e a-steach an luach a chaidh a chuir a-steach don t-slot.

Ma tha ann an aon ghairm kernel gpu_hashtable_insert() tha grunn eileamaidean ann leis an aon iuchair, an uairsin faodar gin de na luachan aca a sgrìobhadh chun phrìomh slot. Thathas den bheachd gu bheil seo àbhaisteach: bidh aon de na prìomh luach a bhios a’ sgrìobhadh tron ​​​​ghairm a’ soirbheachadh, ach leis gu bheil seo uile a’ tachairt ann an co-shìnte taobh a-staigh grunn snàithleanan coileanaidh, chan urrainn dhuinn ro-innse dè an sgrìobhadh cuimhne a bhios mar an tè mu dheireadh.

Sealladh clàr hash

Còd airson iuchraichean lorg:

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

Gus luach iuchair a tha air a stòradh ann an clàr a lorg, bidh sinn ag ath-aithris tron ​​​​raon a’ tòiseachadh le hash na h-iuchrach a tha sinn a’ sireadh. Anns gach slot, bidh sinn a’ dèanamh cinnteach an e an iuchair an tè a tha sinn a’ sireadh, agus ma tha, tillidh sinn a luach. Nì sinn sgrùdadh cuideachd a bheil an iuchair falamh, agus ma tha, sguir sinn den rannsachadh.

Mura h-urrainn dhuinn an iuchair a lorg, tillidh an còd luach falamh.

Faodar na h-obraichean sgrùdaidh sin uile a dhèanamh aig an aon àm tro chuir a-steach agus cuir às. Bidh aon de na ceithir stàitean gu h-àrd aig gach paidhir sa chlàr airson an t-sruth.

Sguab às ann an clàr hash

Còd airson iuchraichean a sguabadh às:

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

Tha cuir às do iuchair air a dhèanamh ann an dòigh neo-àbhaisteach: fàgaidh sinn an iuchair sa chlàr agus comharraich a luach (chan e an iuchair fhèin) mar fhalamh. Tha an còd seo glè choltach ris lookup(), ach a-mhàin nuair a lorgar maids air iuchair, bidh e a’ dèanamh a luach falamh.

Mar a chaidh ainmeachadh gu h-àrd, aon uair 's gu bheil iuchair air a sgrìobhadh gu slot, chan eil e air a ghluasad tuilleadh. Fiù nuair a thèid eileamaid a dhubhadh às bhon chlàr, tha an iuchair fhathast na àite, bidh a luach dìreach falamh. Tha seo a’ ciallachadh nach fheum sinn obrachadh sgrìobhaidh atamach a chleachdadh airson an luach sliotan, oir chan eil e gu diofar a bheil an luach làithreach falamh no nach eil - bidh e fhathast falamh.

Ag ath-mheudachadh clàr hash

Faodaidh tu meud clàr hash atharrachadh le bhith a’ cruthachadh clàr nas motha agus a’ cuir a-steach eileamaidean neo-falamh bhon t-seann chlàr a-steach ann. Cha do chuir mi an gnìomh seo an gnìomh oir bha mi airson an còd sampall a chumail sìmplidh. A bharrachd air an sin, ann am prògraman CUDA, bidh riarachadh cuimhne gu tric air a dhèanamh anns a’ chòd aoigheachd seach ann an kernel CUDA.

Anns an artaigil Clàr hash gun ghlas gun feitheamh ag innse mar a dh’ atharraicheas tu structar dàta a tha fo dhìon glas.

Farpaiseachd

Anns na criomagan còd gnìomh gu h-àrd gpu_hashtable_insert(), _lookup() и _delete() pròiseas aon phaidhir luach-iuchrach aig aon àm. Agus nas ìsle gpu_hashtable_insert(), _lookup() и _delete() pròiseas sreath de chàraidean ann an co-shìnte, gach paidhir ann an snàithlean cur gu bàs GPU air leth:

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

Bidh an clàr hash dìon-glas a’ toirt taic do chuir a-steach, lorg agus cuir às aig an aon àm. Leis gu bheil paidhrichean le prìomh luach an-còmhnaidh ann an aon de cheithir stàitean agus nach eil na h-iuchraichean a’ gluasad, tha an clàr a’ gealltainn ceartachd eadhon nuair a thèid diofar sheòrsaichean obrachaidh a chleachdadh aig an aon àm.

Ach, ma bhios sinn a’ pròiseasadh baidse de chuir-a-steach agus cuir às aig an aon àm, agus ma tha iuchraichean dùblaichte anns an raon cuir a-steach de chàraidean, cha bhith e comasach dhuinn ro-innse dè na paidhrichean a bhios “a’ buannachadh ”- thèid a sgrìobhadh chun chlàr hash mu dheireadh. Canaidh sinn gun tug sinn an còd cuir a-steach le sreath de chàraidean a-steach A/0 B/1 A/2 C/3 A/4. Nuair a bhios an còd a 'crìochnachadh, paidhir B/1 и C/3 gealltainn gum bi iad an làthair air a’ chlàr, ach aig an aon àm nochdaidh gin de na càraidean ann A/0, A/2 no A/4. Faodaidh seo a bhith na dhuilgheadas no nach eil - tha e uile an urra ris an tagradh. Is dòcha gu bheil fios agad ro-làimh nach eil iuchraichean dùblaichte san raon cuir a-steach, no is dòcha nach eil dragh agad dè an luach a chaidh a sgrìobhadh mu dheireadh.

Ma tha seo na dhuilgheadas dhut, feumaidh tu na paidhrichean dùblaichte a sgaradh gu diofar ghairmean siostam CUDA. Ann an CUDA, bidh gnìomhachd sam bith a bhios a’ gairm an kernel an-còmhnaidh a’ crìochnachadh ron ath ghairm kernel (co-dhiù taobh a-staigh aon snàithlean. Ann an diofar snàithleanan, thèid kernels a chuir gu bàs aig an aon àm). Anns an eisimpleir gu h-àrd, ma chanas tu ri aon kernel le A/0 B/1 A/2 C/3, agus am fear eile le A/4, an uairsin an iuchair A gheibh e an luach 4.

A-nis bruidhnidh sinn am bu chòir gnìomhan lookup() и delete() cleachd puing sìmplidh no luaineach gu sreath de chàraidean anns a’ chlàr hash. Sgrìobhainnean CUDA Stàitean a tha:

Faodaidh an neach-cruinneachaidh roghnachadh leughaidhean agus sgrìobhadh gu cuimhne chruinneil no co-roinnte a mheudachadh... Faodar na h-àrdachadh seo a chuir à comas leis a’ phrìomh fhacal volatile: ... tha iomradh sam bith air a’ chaochladair seo air a chur ri chèile ann an stiùireadh fìor leughaidh no sgrìobhaidh cuimhne.

Chan fheum beachdachadh air ceartachd tagradh volatile. Ma chleachdas an t-snàthainn cur gu bàs luach taisgte bho ghnìomhachd a chaidh a leughadh na bu thràithe, bidh e a’ cleachdadh fiosrachadh a tha beagan seann-fhasanta. Ach fhathast, is e seo fiosrachadh bho staid cheart clàr hash aig àm sònraichte den ghairm kernel. Ma dh'fheumas tu am fiosrachadh as ùire a chleachdadh, faodaidh tu an clàr-amais a chleachdadh volatile, ach an uairsin lùghdaichidh an coileanadh beagan: a rèir mo dheuchainnean, nuair a chaidh 32 millean eileamaid a dhubhadh às, chaidh an astar sìos bho 500 millean cuir às / diog gu 450 millean cuir às / diog.

Coileanadh

Anns an deuchainn airson 64 millean eileamaidean a chuir a-steach agus cuir às do 32 millean dhiubh, tha farpais eadar std::unordered_map agus cha mhòr nach eil clàr hash ann airson an GPU:

Clàr hash sìmplidh airson GPU
std::unordered_map chosg e 70 ms a’ cuir a-steach agus a’ toirt air falbh eileamaidean agus an uairsin gan saoradh unordered_map (bidh faighinn cuidhteas milleanan de eileamaidean a’ toirt tòrr ùine, oir a-staigh unordered_map tha ioma-riarachaidhean cuimhne air an dèanamh). A' labhairt gu h-ionraic, std:unordered_map cuingealachaidhean gu tur eadar-dhealaichte. Is e aon snàithlean cur gu bàs CPU a th ’ann, a’ toirt taic do phrìomh luachan de mheud sam bith, a ’coileanadh gu math aig ìrean cleachdaidh àrd, agus a’ nochdadh coileanadh seasmhach às deidh grunn sguabadh às.

B’ e fad a’ bhùird hash airson an GPU agus conaltradh eadar-phrògraman 984 ms. Tha seo a’ toirt a-steach na h-ùine a thathar a’ cur seachad a’ cur a’ bhùird mar chuimhne agus ga dhubhadh às (a’ riarachadh 1 GB de chuimhne aon turas, a bheir beagan ùine ann an CUDA), a’ cuir a-steach agus a’ cuir às do eileamaidean, agus ag aithris thairis orra. Thathas cuideachd a’ toirt aire do gach leth-bhreac gu agus bhon chuimhne cairt bhidio.

Thug an clàr hash fhèin 271 ms airson a chrìochnachadh. Tha seo a’ toirt a-steach an ùine a chaitheas a’ chairt bhidio a’ cuir a-steach agus a’ cuir às do eileamaidean, agus chan eil e a’ toirt aire don ùine a thathar a’ caitheamh a’ dèanamh copaidh dhan chuimhne agus ag aithris thairis air a’ chlàr a thàinig às. Ma tha an clàr GPU beò airson ùine mhòr, no ma tha an clàr hash gu tur mar chuimhneachan air a ’chairt bhidio (mar eisimpleir, gus clàr hash a chruthachadh a thèid a chleachdadh le còd GPU eile agus chan e am pròiseasar meadhanach), an uairsin tha toradh na deuchainn buntainneach.

Tha an clàr hash airson cairt bhidio a’ nochdadh àrd-choileanadh mar thoradh air toradh àrd agus co-shìnteadh gnìomhach.

uireasbhaidhean

Tha beagan chùisean aig ailtireachd clàr hash airson a bhith mothachail:

  • Tha sgrùdadh sreathach air a bhacadh le bhith a’ cruinneachadh, a tha ag adhbhrachadh nach bi na h-iuchraichean sa chlàr air an cur nas lugha na gu foirfe.
  • Chan eil iuchraichean air an toirt air falbh leis a’ ghnìomh delete agus thar ùine tha iad a' bearradh am bòrd.

Mar thoradh air an sin, faodaidh coileanadh clàr hash a dhol sìos mean air mhean, gu sònraichte ma tha e ann airson ùine mhòr agus gu bheil grunn chuir a-steach agus cuir às. Is e aon dòigh air na h-eas-bhuannachdan sin a lughdachadh a bhith ag ath-shuidheachadh a-steach do bhòrd ùr le ìre cleachdaidh gu math ìosal agus a’ sìoladh a-mach na h-iuchraichean a chaidh a thoirt air falbh aig àm an ath-shuidheachadh.

Gus na cùisean a chaidh a mhìneachadh a nochdadh, cleachdaidh mi an còd gu h-àrd gus clàr a chruthachadh le 128 millean eileamaid agus lùb tro 4 millean eileamaid gus am bi mi air 124 millean sliotan a lìonadh (ìre cleachdaidh timcheall air 0,96). Seo an clàr toraidh, is e gairm kernel CUDA a th’ anns gach sreath gus 4 millean eileamaid ùr a chuir a-steach do aon chlàr hash:

Ìre cleachdaidh
Faid cuir a-steach 4 eileamaidean

0,00
11,608448 ms (361,314798 millean iuchraichean / diog.)

0,03
11,751424 ms (356,918799 millean iuchraichean / diog.)

0,06
11,942592 ms (351,205515 millean iuchraichean / diog.)

0,09
12,081120 ms (347,178429 millean iuchraichean / diog.)

0,12
12,242560 ms (342,600233 millean iuchraichean / diog.)

0,16
12,396448 ms (338,347235 millean iuchraichean / diog.)

0,19
12,533024 ms (334,660176 millean iuchraichean / diog.)

0,22
12,703328 ms (330,173626 millean iuchraichean / diog.)

0,25
12,884512 ms (325,530693 millean iuchraichean / diog.)

0,28
13,033472 ms (321,810182 millean iuchraichean / diog.)

0,31
13,239296 ms (316,807174 millean iuchraichean / diog.)

0,34
13,392448 ms (313,184256 millean iuchraichean / diog.)

0,37
13,624000 ms (307,861434 millean iuchraichean / diog.)

0,41
13,875520 ms (302,280855 millean iuchraichean / diog.)

0,44
14,126528 ms (296,909756 millean iuchraichean / diog.)

0,47
14,399328 ms (291,284699 millean iuchraichean / diog.)

0,50
14,690304 ms (285,515123 millean iuchraichean / diog.)

0,53
15,039136 ms (278,892623 millean iuchraichean / diog.)

0,56
15,478656 ms (270,973402 millean iuchraichean / diog.)

0,59
15,985664 ms (262,379092 millean iuchraichean / diog.)

0,62
16,668673 ms (251,627968 millean iuchraichean / diog.)

0,66
17,587200 ms (238,486174 millean iuchraichean / diog.)

0,69
18,690048 ms (224,413765 millean iuchraichean / diog.)

0,72
20,278816 ms (206,831789 millean iuchraichean / diog.)

0,75
22,545408 ms (186,038058 millean iuchraichean / diog.)

0,78
26,053312 ms (160,989275 millean iuchraichean / diog.)

0,81
31,895008 ms (131,503463 millean iuchraichean / diog.)

0,84
42,103294 ms (99,619378 millean iuchraichean / diog.)

0,87
61,849056 ms (67,815164 millean iuchraichean / diog.)

0,90
105,695999 ms (39,682713 millean iuchraichean / diog.)

0,94
240,204636 ms (17,461378 millean iuchraichean / diog.)

Mar a bhios cleachdadh a’ dol am meud, bidh coileanadh a’ dol sìos. Chan eil seo ion-mhiannaichte sa mhòr-chuid de chùisean. Ma chuireas tagradh eileamaidean a-steach do bhòrd agus an uairsin gan tilgeadh air falbh (mar eisimpleir, nuair a thathar a’ cunntadh fhaclan ann an leabhar), chan eil seo na dhuilgheadas. Ach ma chleachdas an aplacaid clàr hash fad-ùine (mar eisimpleir, ann an deasaiche grafaigs gus pàirtean neo-fholamh de dhealbhan a stòradh far am bi an neach-cleachdaidh gu tric a’ cuir a-steach agus a’ cuir às do fhiosrachadh), faodaidh an giùlan seo a bhith na dhuilgheadas.

Agus thomhais e doimhneachd sgrùdaidh clàr hash às deidh 64 millean cuir a-steach (factar cleachdaidh 0,5). B 'e an doimhneachd cuibheasach 0,4774, agus mar sin bha a' mhòr-chuid de na h-iuchraichean an dàrna cuid anns an t-slot as fheàrr no aon slot air falbh bhon t-suidheachadh as fheàrr. B' e an cuibhreann stoc as lugha de 60.

An uairsin thomhais mi an doimhneachd sgrùdaidh air bòrd le 124 millean cuir a-steach (factar cleachdaidh 0,97). Bha an doimhneachd cuibheasach mu thràth 10,1757, agus an ìre as àirde - 6474 (!!). Bidh coileanadh mothachaidh sreathach a’ tuiteam gu mòr aig ìrean cleachdaidh àrd.

Tha e nas fheàrr ìre cleachdaidh a’ bhùird hash seo a chumail ìosal. Ach an uairsin bidh sinn ag àrdachadh coileanadh aig cosgais caitheamh cuimhne. Gu fortanach, a thaobh iuchraichean agus luachan 32-bit, faodar seo a dhearbhadh. Ma chumas sinn san eisimpleir gu h-àrd, ann an clàr le 128 millean eileamaidean, am feart cleachdaidh de 0,25, chan urrainn dhuinn barrachd air 32 millean eileamaid a chuir ann, agus thèid na 96 millean sliotan a tha air fhàgail a chall - 8 bytes airson gach paidhir. , 768 MB de chuimhne caillte.

Thoir an aire gu bheil sinn a’ bruidhinn mu dheidhinn call cuimhne cairt bhidio, a tha na ghoireas nas luachmhoire na cuimhne siostam. Ged a tha co-dhiù 4 GB de chuimhne aig a ’mhòr-chuid de chairtean grafaiceachd deasg ùr-nodha a bheir taic do CUDA (aig àm sgrìobhaidh, tha 2080 GB aig an NVIDIA 11 Ti), cha bhiodh e fhathast mar an co-dhùnadh as glic na suimean sin a chall.

Nas fhaide air adhart sgrìobhaidh mi barrachd mu bhith a’ cruthachadh clàran hash airson cairtean bhidio aig nach eil duilgheadasan le doimhneachd sgrùdaidh, a bharrachd air dòighean air sliotan a chaidh a dhubhadh às ath-chleachdadh.

Tomhas doimhneachd fuaim

Gus doimhneachd sgrùdaidh iuchrach a dhearbhadh, is urrainn dhuinn hash na h-iuchrach (an clàr-amais air leth freagarrach) a thoirt a-mach às a’ chlàr clàr fhèin:

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

Air sgàth draoidheachd àireamhan dà-chànanach dà dhà agus leis gu bheil comas a’ bhùird hash dà ri cumhachd dhà, obraichidh an dòigh-obrach seo eadhon nuair a thèid am prìomh chlàr-amais a ghluasad gu toiseach a’ chlàir. Gabhaidh sinn iuchair a tha hashed gu 1, ach air a chuir a-steach gu slot 3. An uairsin airson bòrd le comas 4 gheibh sinn (3 — 1) & 3, a tha co-ionann ri 2.

co-dhùnadh

Ma tha ceistean no beachdan agad, cuir post-d thugam aig Twitter no fosgail cuspair ùr a-staigh tasgaidh.

Chaidh an còd seo a sgrìobhadh fo bhrosnachadh bho artaigilean sàr-mhath:

San àm ri teachd, cumaidh mi a’ sgrìobhadh mu bhuileachadh clàr hash airson cairtean bhidio agus a’ dèanamh anailis air an coileanadh. Tha na planaichean agam a’ toirt a-steach slabhraidh, hashing Robin Hood, agus hashing cuthaige a’ cleachdadh gnìomhachd atamach ann an structaran dàta a tha càirdeil do GPU.

Source: www.habr.com

Cuir beachd ann