Na buga shi akan Github .
Teburin hash ne mai sauƙi na GPU mai ikon sarrafa ɗaruruwan miliyoyin abubuwan sakawa a sakan daya. A kan kwamfutar tafi-da-gidanka na NVIDIA GTX 1060, lambar ta sanya nau'ikan ƙima-ƙimar miliyan 64 da ka ƙirƙira a cikin kusan ms 210 kuma tana cire nau'i-nau'i miliyan 32 a cikin kusan 64 ms.
Wato, gudun kan kwamfutar tafi-da-gidanka ya kai kusan 300 miliyan shigarwa / s da 500 miliyan sharewa / sec.
An rubuta teburin a cikin CUDA, kodayake ana iya amfani da wannan fasaha ga HLSL ko GLSL. Aiwatar yana da iyakoki da yawa don tabbatar da babban aiki akan katin bidiyo:
- Maɓallai 32-bit kawai da ƙimar iri ɗaya ake sarrafa su.
- Teburin zanta yana da ƙayyadaddun girma.
- Kuma wannan girman dole ne ya zama daidai da biyu zuwa ikon.
Don maɓalli da ƙima, kuna buƙatar tanadin alamar ƙima mai sauƙi (a cikin lambar da ke sama wannan shine 0xffffffff).
Teburin Hash ba tare da makullai ba
Teburin zanta yana amfani da buɗaɗɗen adireshi tare da , wato, shi ne kawai tsararrun maɓalli-darajar nau'i-nau'i waɗanda aka adana a cikin ƙwaƙwalwar ajiya kuma suna da kyakkyawan aikin cache. Ba za a iya faɗi iri ɗaya ba don sarƙaƙƙiya, wanda ya haɗa da neman mai nuni a cikin jerin da aka haɗa. Teburin zanta shine mai sauƙi tsararrun adana abubuwa KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Girman tebur yana da iko na biyu, ba lambar farko ba, saboda umarni ɗaya mai sauri ya isa a yi amfani da mashin pow2/AND, amma ma'aikacin modules yana da hankali sosai. Wannan yana da mahimmanci a yanayin binciken linzamin kwamfuta, tunda a cikin duban tebur na linzamin kwamfuta dole ne a nannade fihirisar ramin a kowane ramin. Kuma a sakamakon haka, ana ƙara farashin aikin modulo a kowane rami.
Tebur kawai yana adana maɓalli da ƙima ga kowane kashi, ba hash na maɓallin ba. Tun da tebur kawai yana adana maɓallan 32-bit, ana ƙididdige zanta da sauri. Lambar da ke sama tana amfani da hash na Murmur3, wanda ke yin ƴan sauye-sauye, XORs da ninkawa.
Teburin zanta yana amfani da dabarun kariyar kullewa waɗanda ba su da tsarin tsarin ƙwaƙwalwa. Ko da wasu rubuta ayyukan sun ɓata tsarin wasu irin waɗannan ayyuka, tebur ɗin zanta zai ci gaba da kiyaye yanayin daidai. Za mu yi magana game da wannan a kasa. Dabarar tana aiki da kyau tare da katunan bidiyo waɗanda ke tafiyar da dubban zaren lokaci guda.
Maɓallai da ƙididdiga a cikin tebur ɗin hash an fara su zuwa komai.
Ana iya canza lambar don sarrafa maɓallan 64-bit da ƙimar kuma. Maɓallai suna buƙatar karantawa, rubutawa, da kwatanta-da-swap ayyukan. Kuma dabi'u suna buƙatar karantawa da rubuta ayyukan atomic. An yi sa'a, a cikin CUDA, ayyukan karanta-rubutu don ƙimar 32- da 64-bit sune atomic muddin suna daidaitawa ta zahiri (duba ƙasa). ), da katunan bidiyo na zamani suna goyan bayan 64-bit atomic kwatanta-da-musanyawa ayyuka. Tabbas, lokacin motsawa zuwa 64 bits, aikin zai ragu kaɗan.
Yanayin tebur Hash
Kowane maɓalli-darajar maɓalli a cikin tebur ɗin zanta na iya samun ɗayan jihohi huɗu:
- Maɓalli da ƙima babu komai. A wannan yanayin, an fara fara zanta tebur.
- An rubuta maɓalli, amma har yanzu ba a rubuta ƙimar ba. Idan wani zaren yana karanta bayanai a halin yanzu, sai ya dawo fanko. Wannan shi ne al'ada, irin wannan abu zai faru idan wani zaren kisa ya yi aiki kadan a baya, kuma muna magana ne game da tsarin bayanan lokaci guda.
- Duka maɓalli da ƙimar ana rubuta su.
- Ƙimar tana samuwa ga sauran zaren aiwatarwa, amma maɓallin bai riga ya ƙare ba. Wannan na iya faruwa saboda ƙirar shirye-shiryen CUDA yana da ƙirar ƙwaƙwalwar ajiya mara izini. Wannan al'ada ce; a kowane yanayi, maɓalli har yanzu fanko ne, ko da ƙimar ba haka take ba.
Wani muhimmin nuance shi ne cewa da zarar an rubuta maɓallin zuwa ramin, ba zai sake motsawa ba - ko da maɓallin ya share, za mu yi magana game da wannan a ƙasa.
Lambar tebur ɗin zanta har ma yana aiki tare da ƙirar ƙwaƙwalwar ajiya da aka ba da oda wanda ba a san tsarin da ake karantawa da rubuta ƙwaƙwalwar ajiya ba. Yayin da muke duban sakawa, dubawa, da gogewa a cikin tebur ɗin zanta, tuna cewa kowane maɓalli-daraja biyu suna cikin ɗaya daga cikin jihohi huɗu da aka kwatanta a sama.
Saka cikin tebur ɗin zanta
Aikin CUDA wanda ke shigar da nau'i-nau'i-daraja maɓalli a cikin tebur ɗin hash yayi kama da haka:
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);
}
}
Don saka maɓalli, lambar tana ƙara ta cikin jeri na tebur ɗin zanta yana farawa da zaton maɓallin da aka saka. Kowane rami a cikin tsararru yana yin aikin kwatanta-da-swap na atomic wanda ke kwatanta maɓalli a cikin wannan ramin zuwa komai. Idan an gano rashin daidaituwa, ana sabunta maɓalli a cikin ramin tare da maɓallin da aka saka, sannan a dawo da maɓallin ramin na asali. Idan wannan maɓalli na asali ba komai bane ko kuma yayi daidai da maɓallin da aka saka, to lambar ta sami ramin da ya dace don sakawa kuma a saka ƙimar da aka saka a cikin ramin.
Idan a cikin kira kernel daya gpu_hashtable_insert() akwai abubuwa da yawa tare da maɓalli iri ɗaya, sannan kowane ɗayan ƙimar su ana iya rubuta shi zuwa maɓallin maɓalli. Ana ɗaukar wannan al'ada: ɗaya daga cikin maɓalli-darajar da aka rubuta yayin kiran zai yi nasara, amma tunda duk wannan yana faruwa a layi daya tsakanin zaren aiwatarwa da yawa, ba za mu iya hango ko wane rubutun ƙwaƙwalwar zai zama na ƙarshe ba.
Duba tebur Hash
Lambar don maɓallan bincike:
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);
}
}
Don nemo darajar maɓalli da aka adana a cikin tebur, muna ƙididdigewa ta hanyar tsararrun farawa da zanta na maɓallin da muke nema. A kowane ramin, muna bincika idan maɓalli shine wanda muke nema, kuma idan haka ne, zamu dawo da ƙimarsa. Muna kuma bincika idan maɓalli ba komai, kuma idan haka ne, muna soke binciken.
Idan ba za mu iya nemo maɓalli ba, lambar zata dawo da fanko darajar.
Duk waɗannan ayyukan bincike ana iya yin su a lokaci guda ta hanyar sakawa da gogewa. Kowane nau'i-nau'i a cikin tebur za su sami ɗaya daga cikin jihohi huɗu da aka kwatanta a sama don kwarara.
Sharewa a cikin tebur ɗin zanta
Lambar don share maɓalli:
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);
}
}
Ana yin share maɓalli ta hanyar da ba a saba gani ba: muna barin maɓallin a cikin tebur kuma mu sanya alamar darajarsa (ba maɓalli da kansa ba) a matsayin fanko. Wannan code yayi kama da lookup(), sai dai idan aka sami ashana akan maɓalli, sai ya mayar da ƙimarsa fanko.
Kamar yadda aka ambata a sama, da zarar an rubuta maɓalli zuwa ramin, ba a sake motsa shi. Ko da an goge wani abu daga tebur, maɓallin yana nan a wurin, ƙimarsa kawai ta zama fanko. Wannan yana nufin cewa ba ma buƙatar yin amfani da aikin rubuta atomic don ƙimar ramin, saboda ba kome ko darajar yanzu ba ta da komai ko a'a - har yanzu za ta zama fanko.
Ana sake fasalin teburin zanta
Kuna iya canza girman teburin zanta ta hanyar ƙirƙirar tebur mafi girma da saka abubuwan da ba su da komai daga tsohon tebur a ciki. Ban aiwatar da wannan aikin ba saboda ina so in kiyaye lambar samfurin mai sauƙi. Bugu da ƙari, a cikin shirye-shiryen CUDA, ana yin rarraba ƙwaƙwalwar ajiya sau da yawa a cikin lambar mai watsa shiri maimakon a cikin CUDA kernel.
Labarin ya bayyana yadda ake gyara irin wannan tsarin bayanan da aka karewa.
Gasa
A cikin snippets code na sama aiki gpu_hashtable_insert(), _lookup() и _delete() aiwatar da maɓalli-daraja guda biyu a lokaci guda. Kuma ƙasa gpu_hashtable_insert(), _lookup() и _delete() aiwatar da tsararrun nau'i-nau'i a layi daya, kowane nau'i-nau'i a cikin zaren aiwatar da GPU daban:
// 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);
}
}
Teburin zanta mai jure wa kulle yana goyan bayan abubuwan sawa lokaci guda, dubawa, da gogewa. Saboda maɓalli-darajar nau'i-nau'i koyaushe suna cikin ɗaya daga cikin jihohi huɗu kuma maɓallan ba sa motsawa, tebur yana ba da garantin daidai koda lokacin da ake amfani da nau'ikan ayyuka daban-daban a lokaci guda.
Koyaya, idan muka aiwatar da juzu'in shigarwa da gogewa a layi daya, kuma idan tsarin shigar da nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau`ikan da aka shigar da su shigar da su sun ƙunshi maɓallan maɓalli, to ba za mu iya yin hasashen waɗanne nau'ikan nau'ikan "nasara" za su yi nasara ba - za a rubuta su zuwa teburin zanta a ƙarshe. Bari mu ce mun kira lambar shigarwa tare da tsararrun shigarwar nau'i-nau'i A/0 B/1 A/2 C/3 A/4. Lokacin da lambar ta cika, nau'i-nau'i B/1 и C/3 an ba da tabbacin kasancewa a cikin tebur, amma a lokaci guda kowane ɗayan nau'i-nau'i zai bayyana a ciki A/0, A/2 ko A/4. Wannan yana iya ko bazai zama matsala ba - duk ya dogara da aikace-aikacen. Kila ka sani a gaba cewa babu kwafin maɓallai a cikin tsararrun shigarwar, ko kuma ba za ka damu da wace ƙima aka rubuta ta ƙarshe ba.
Idan wannan matsala ce a gare ku, to kuna buƙatar raba nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan nau'ikan kiran tsarin CUDA daban-daban. A cikin CUDA, duk wani aiki da ke kiran kernel koyaushe yana cika kafin kiran kernel na gaba (aƙalla cikin zare ɗaya. A cikin zaren daban-daban, ana aiwatar da kernel a layi daya). A cikin misalin da ke sama, idan kun kira kwaya ɗaya da A/0 B/1 A/2 C/3, da sauran tare da A/4, sannan kullin A zai samu darajar 4.
Yanzu bari muyi magana game da ko ayyuka ya kamata lookup() и delete() yi amfani da ma'ana a sarari ko mai canzawa zuwa tsararrun nau'i-nau'i a cikin tebur ɗin zanta. Yana cewa:
Mai tarawa na iya zaɓar inganta karantawa da rubutu zuwa ƙwaƙwalwar ajiya ta duniya ko rabawa… Ana iya kashe waɗannan haɓakawa ta amfani da kalmar maɓalli.
volatile: ... duk wani nuni ga wannan mabambanta an haɗa shi cikin ainihin ƙwaƙwalwar karantawa ko rubuta umarni.
La'akarin daidaito baya buƙatar aikace-aikace volatile. Idan zaren kisa yana amfani da ƙimar da aka adana daga aikin karantawa a baya, to zai kasance yana amfani da bayanan da suka wuce. Amma duk da haka, wannan bayani ne daga daidai yanayin tebur ɗin zanta a wani lokaci na kiran kernel. Idan kuna buƙatar amfani da sabbin bayanai, zaku iya amfani da fihirisar volatile, amma sai aikin zai ragu kaɗan: bisa ga gwaje-gwaje na, lokacin da ake share abubuwa miliyan 32, saurin ya ragu daga sharewa miliyan 500 / sec zuwa 450 deletions / sec.
Yawan aiki
A cikin gwajin saka abubuwa miliyan 64 da kuma share miliyan 32 daga cikinsu, gasa tsakanin std::unordered_map kuma kusan babu teburin zanta don GPU:
std::unordered_map ya kashe 70 ms wajen sakawa da cire abubuwa sannan a 'yantar da su unordered_map (cire miliyoyin abubuwa yana ɗaukar lokaci mai yawa, saboda ciki unordered_map an yi rabon adadin ƙwaƙwalwar ajiya da yawa). Maganar gaskiya, std:unordered_map mabanbanta hani. Zaren aiwatar da CPU guda ɗaya ne, yana goyan bayan mahimman ƙimar kowane girman, yana aiki da kyau a ƙimar amfani mai yawa, kuma yana nuna ingantaccen aiki bayan gogewa da yawa.
Tsawon lokacin teburin zanta don GPU da sadarwa tsakanin shirye-shirye shine 984 ms. Wannan ya hada da lokacin da aka kashe wajen ajiye tebur a cikin ƙwaƙwalwar ajiya da goge shi (waɗanda 1 GB na ƙwaƙwalwar ajiya lokaci ɗaya, wanda ke ɗaukar ɗan lokaci a cikin CUDA), sakawa da goge abubuwa, da sake maimaita su. Ana la'akari da duk kwafi zuwa ko daga ƙwaƙwalwar ajiyar katin bidiyo.
Teburin zanta da kansa ya ɗauki 271 ms don kammalawa. Wannan ya haɗa da lokacin da katin bidiyo ya kashe don sakawa da share abubuwa, kuma baya la'akari da lokacin da aka kashe don yin kwafi zuwa ƙwaƙwalwar ajiya da maimaita kan teburin da aka samu. Idan tebur na GPU yana rayuwa na dogon lokaci, ko kuma idan tebur ɗin hash yana ƙunshe gaba ɗaya a cikin ƙwaƙwalwar katin bidiyo (misali, don ƙirƙirar tebur ɗin zanta wanda sauran lambar GPU za ta yi amfani da shi ba na tsakiya ba), to. sakamakon gwajin ya dace.
Teburin zanta don katin bidiyo yana nuna babban aiki saboda babban kayan aiki da daidaitawa mai aiki.
shortcomings
Hash table architecture yana da ƴan al'amura da yakamata ku sani:
- Binciken layi yana da cikas ta hanyar tari, wanda ke haifar da maɓallai a cikin tebur don sanya ƙasa da daidai.
- Ba a cire maɓallai ta amfani da aikin
deletekuma a kan lokaci suna cika tebur.
Sakamakon haka, aikin tebur ɗin zanta na iya raguwa sannu a hankali, musamman idan ya daɗe yana da abubuwan sakawa da gogewa. Hanya ɗaya don rage waɗannan lahani ita ce sake kunna sabon tebur tare da ƙarancin amfani da ƙarancin amfani da kuma tace maɓallan da aka cire yayin sakewa.
Don kwatanta batutuwan da aka bayyana, zan yi amfani da lambar da ke sama don ƙirƙirar tebur mai abubuwa miliyan 128 da madauki ta hanyar abubuwa miliyan 4 har sai na cika ramummuka miliyan 124 (yawan amfani da kusan 0,96). Anan ga teburin sakamako, kowane layi shine kiran kwaya na CUDA don saka sabbin abubuwa miliyan 4 cikin tebur hash ɗaya:
Yawan amfani
Tsawon lokacin shigar abubuwa 4
0,00
11,608448 ms (361,314798 maɓallai miliyan/sek.)
0,03
11,751424 ms (356,918799 maɓallai miliyan/sek.)
0,06
11,942592 ms (351,205515 maɓallai miliyan/sek.)
0,09
12,081120 ms (347,178429 maɓallai miliyan/sek.)
0,12
12,242560 ms (342,600233 maɓallai miliyan/sek.)
0,16
12,396448 ms (338,347235 maɓallai miliyan/sek.)
0,19
12,533024 ms (334,660176 maɓallai miliyan/sek.)
0,22
12,703328 ms (330,173626 maɓallai miliyan/sek.)
0,25
12,884512 ms (325,530693 maɓallai miliyan/sek.)
0,28
13,033472 ms (321,810182 maɓallai miliyan/sek.)
0,31
13,239296 ms (316,807174 maɓallai miliyan/sek.)
0,34
13,392448 ms (313,184256 maɓallai miliyan/sek.)
0,37
13,624000 ms (307,861434 maɓallai miliyan/sek.)
0,41
13,875520 ms (302,280855 maɓallai miliyan/sek.)
0,44
14,126528 ms (296,909756 maɓallai miliyan/sek.)
0,47
14,399328 ms (291,284699 maɓallai miliyan/sek.)
0,50
14,690304 ms (285,515123 maɓallai miliyan/sek.)
0,53
15,039136 ms (278,892623 maɓallai miliyan/sek.)
0,56
15,478656 ms (270,973402 maɓallai miliyan/sek.)
0,59
15,985664 ms (262,379092 maɓallai miliyan/sek.)
0,62
16,668673 ms (251,627968 maɓallai miliyan/sek.)
0,66
17,587200 ms (238,486174 maɓallai miliyan/sek.)
0,69
18,690048 ms (224,413765 maɓallai miliyan/sek.)
0,72
20,278816 ms (206,831789 maɓallai miliyan/sek.)
0,75
22,545408 ms (186,038058 maɓallai miliyan/sek.)
0,78
26,053312 ms (160,989275 maɓallai miliyan/sek.)
0,81
31,895008 ms (131,503463 maɓallai miliyan/sek.)
0,84
42,103294 ms (99,619378 maɓallai miliyan/sek.)
0,87
61,849056 ms (67,815164 maɓallai miliyan/sek.)
0,90
105,695999 ms (39,682713 maɓallai miliyan/sek.)
0,94
240,204636 ms (17,461378 maɓallai miliyan/sek.)
Yayin da amfani ke ƙaruwa, aikin yana raguwa. Wannan ba kyawawa bane a mafi yawan lokuta. Idan aikace-aikacen ya saka abubuwa a cikin tebur sannan ya watsar da su (misali, lokacin ƙidayar kalmomi a cikin littafi), to wannan ba matsala ba ce. Amma idan aikace-aikacen yana amfani da tebur na zanta na dogon lokaci (misali, a cikin editan zane don adana sassan hotuna marasa fanko inda mai amfani ke yawan sakawa da share bayanai), to wannan hali na iya zama matsala.
Kuma auna zurfin binciken tebur ɗin zanta bayan an shigar da miliyan 64 (fasali mai amfani 0,5). Matsakaicin zurfin ya kasance 0,4774, don haka yawancin maɓallai sun kasance ko dai a cikin mafi kyawun yuwuwar ramin ko rami ɗaya nesa da mafi kyawun matsayi. Matsakaicin zurfin sauti shine 60.
Daga nan na auna zurfin bincike akan tebur tare da abubuwan sakawa miliyan 124 (fasali mai amfani 0,97). Matsakaicin zurfin ya riga ya kasance 10,1757, kuma matsakaicin - 6474 (!!). Ayyukan ji na layi yana raguwa sosai a ƙimar amfani mai yawa.
Zai fi kyau a kiyaye ƙimar amfani da wannan tebur ɗin hash ƙasa kaɗan. Amma sai mu ƙara yawan aiki a kuɗin amfani da ƙwaƙwalwar ajiya. Abin farin ciki, a cikin yanayin maɓallan 32-bit da ƙima, wannan na iya zama barata. Idan a cikin misalin da ke sama, a cikin tebur tare da abubuwa miliyan 128, muna kiyaye ƙimar amfani na 0,25, to, ba za mu iya sanya abubuwa sama da miliyan 32 a ciki ba, kuma sauran ramummuka miliyan 96 za su ɓace - 8 bytes ga kowane ɗayan biyu. , 768 MB na ƙwaƙwalwar ajiya.
Lura cewa muna magana ne game da asarar ƙwaƙwalwar ajiyar katin bidiyo, wanda shine mafi mahimmancin albarkatu fiye da ƙwaƙwalwar tsarin. Kodayake yawancin katunan zane na zamani waɗanda ke goyan bayan CUDA suna da aƙalla 4 GB na ƙwaƙwalwar ajiya (a lokacin rubutawa, NVIDIA 2080 Ti yana da 11 GB), har yanzu ba zai zama yanke shawara mafi hikima don rasa irin wannan adadin ba.
Daga baya zan rubuta ƙarin game da ƙirƙirar tebur na zanta don katunan bidiyo waɗanda ba su da matsala tare da zurfin bincike, da kuma hanyoyin sake amfani da ramummuka da aka goge.
Ma'aunin zurfin sauti
Don tantance zurfin bincike na maɓalli, za mu iya fitar da zanta na maɓalli (madaidaicin tebur ɗin sa) daga ainihin ma'aunin tebur ɗinsa:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Saboda sihirin lambobi biyu masu dacewa da lambobi biyu da kuma gaskiyar cewa ƙarfin teburin zanta ya kasance biyu zuwa ƙarfin biyu, wannan tsarin zai yi aiki ko da lokacin da aka motsa maɓallin maɓallin zuwa farkon tebur. Bari mu ɗauki maɓalli wanda ya haɗe zuwa 1, amma an saka shi a cikin Ramin 3. Sannan ga tebur mai ƙarfi 4 muna samun (3 — 1) & 3, wanda yayi daidai da 2.
ƙarshe
Idan kuna da tambayoyi ko sharhi, da fatan za a yi mini imel a ko bude wani sabon batu a ciki.
An rubuta wannan lambar a ƙarƙashin wahayi daga ingantattun labarai:
A nan gaba, zan ci gaba da rubuta game da aiwatar da tebur na hash don katunan bidiyo da nazarin ayyukan su. Shirye-shiryena sun haɗa da sarƙaƙƙiya, Robin Hood hashing, da cuckoo hashing ta amfani da ayyukan atomic a cikin tsarin bayanan da ke abokantaka na GPU.
source: www.habr.com