I whakairihia e ahau ki runga i a Github .
He tepu hash GPU ngawari e taea ana ki te tukatuka i nga rau miriona whakaurunga mo ia hekona. I runga i taku pona NVIDIA GTX 1060, ka whakauruhia e te waehere te 64 miriona i hangaia matapōkeretia nga takirua uara-matua i roto i te 210 ms me te tango i te 32 miriona takirua i roto i te 64 ms.
Arā, ko te tere i runga i te pona he tata ki te 300 miriona whakauru / hekona me te 500 miriona whakakore / hekona.
Kua tuhia te ripanga ki te CUDA, ahakoa ka taea te whakamahi i taua tikanga ki te HLSL, ki te GLSL ranei. He maha nga here o te whakatinanatanga hei whakarite i nga mahi teitei i runga i te kaari ataata:
- Ko nga taviri 32-bit anake me nga uara ano ka tukatukahia.
- He rahi te rahi o te ripanga hash.
- A ko tenei rahi me rite ki te rua ki te kaha.
Mo nga taviri me nga uara, me rahui koe i tetahi tohu wehewehe ngawari (i te waehere o runga ake nei ko 0xffffffff).
Tepu Hash kahore raka
Ka whakamahia e te ripanga hash te korero tuwhera me te , ara, he huinga takirua uara-matua ka rongoa ki roto i te mahara, he pai ake te mahi keteroki. Kaore e taea te korero mo te mekameka, me rapu he tohu i roto i te rarangi hono. Ko te ripanga hash he huānga rokiroki ngawari KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Ko te rahi o te ripanga ko te mana o te rua, ehara i te tau matua, no te mea he nui te ako tere mo te tono i te kanohi pow2/AND, engari he puhoi ake te kaiwhakahaere modulus. He mea nui tenei i roto i te take o te tirotiro rarangi, i te mea i roto i te titiro tepu rarangi me takai te tohu mokamoka ki ia mokamoka. Na ko te mutunga, ko te utu o te mahi ka taapirihia te modulo ki ia mokamoka.
Pupuri ana te ripanga i te kī me te uara mō ia huānga, ehara i te hash o te kī. I te mea ko te ripanga e pupuri ana i nga taviri moka-32, ka tere te tatau i te hash. Ko te waehere i runga ake nei e whakamahi ana i te Hash Murmur3, he iti noa nga nekehanga, XOR me nga whakareatanga.
Ka whakamahia e te ripanga hash nga tikanga whakamarumaru maukati e motuhake ana i te raupapa mahara. Ahakoa ka whakararuraru etahi o nga mahi tuhi i te ota o etahi atu mahi pera, ka mau tonu te ripanga hash i te ahua tika. Ka korero tatou mo tenei i raro nei. He pai te mahi me nga kaari ataata e rere ana nga mano o nga miro i te wa kotahi.
Ko nga taviri me nga uara o te ripanga hash ka arawhiti ki te putua.
Ka taea te whakarereke i te waehere hei hapai i nga taviri 64-bit me nga uara. Ko nga taviri e hiahia ana ki te panui ngota, te tuhi, me nga mahi whakataurite me te whakawhiti. A ko nga uara e hiahia ana ki nga mahi panui me te tuhi ngota. Waimarie, i roto i te CUDA, ko nga mahi panui-tuhi mo nga uara 32- me te 64-bit he ngota i te mea he rite tonu te ahua (tirohia ki raro). ), me nga kaari ataata hou e tautoko ana i nga mahi whakataurite-me-whakawhitinga ngota 64-bit. Ae ra, ka neke ki te 64 bits, ka iti haere te mahi.
Te ahua ripanga Hash
Ka taea e ia takirua uara-matua o te ripanga hash tetahi o nga ahuatanga e wha:
- Kei te putua te matua me te uara. I tenei ahua, ka arawhiti te ripanga hash.
- Kua tuhia te matua, engari kaore ano kia tuhia te uara. Mena kei te panui raraunga tetahi atu miro, ka hoki kau mai. He mea noa tenei, he rite ano te ahua mena kua mahi tetahi atu miro o te mahi i mua ake nei, a kei te korero tatou mo te hanganga raraunga kotahi.
- Ka tuhia te matua me te uara.
- Kei te waatea te uara ki etahi atu miro o te mahi, engari kaore ano te matua. Ka tupu pea tenei na te mea he tauira mahara kore te tauira o te CUDA. He mea noa tenei; ahakoa he aha, kei te noho kau tonu te ki, ahakoa kua kore te uara.
Ko tetahi ahuatanga nui ko te mea kua tuhia te ki ki te mokamoka, kua kore e neke - ahakoa kua mukua te ki, ka korero tatou mo tenei i raro nei.
Ko te waehere ripanga hash ka mahi tahi me nga tauira mahara kua ota kare i te mohiotia te raupapa o te panui me te tuhi. I a tatou e titiro ana ki te whakauru, te rapu, me te muku i roto i te ripanga hash, maharatia kei roto ia takirua uara-matua i tetahi o nga ahuatanga e wha kua whakaahuahia i runga ake nei.
Te whakauru ki roto i te ripanga hash
Ko te mahi CUDA e whakauru ana i nga takirua uara-matua ki roto i te ripanga hash he penei te ahua:
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);
}
}
Hei kōkuhu kī, ka huri te waehere i roto i te huinga ripanga hash ka timata ki te hash o te kī whakauru. Ka mahia e ia mokamoka i roto i te huinga he mahi whakatairite-a-whakawhitinga ngota e whakataurite ana i te ki o taua mokamoka ki te putua. Mena ka kitea he koretake, ka whakahōuhia te kī i roto i te mokamoka ki te kī kua whakauruhia, ka whakahokia te kī mokamoka taketake. Mēnā he putua tēnei kī taketake, he ōrite rānei ki te kī i kōkuhuhia, ka kitea e te waehere he mokamoka pai mo te whakauru me te whakauru i te uara kua whakauruhia ki roto i te mokamoka.
Mena i roto i te karanga kernel kotahi gpu_hashtable_insert() he maha nga huānga me te kī kotahi, katahi ka taea te tuhi tetahi o o raatau uara ki te mokamoka matua. He mea noa tenei: ka angitu tetahi o nga uara-matua tuhi i te wa o te waeatanga, engari i te mea ka puta whakarara enei katoa i roto i te maha o nga miro o te mahi, kaore e taea e taatau te matapae ko te tuhi mahara te mea whakamutunga.
Tirohanga ripanga Hash
Waehere mo te rapu taviri:
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);
}
}
Ki te kimi i te uara o te taviri kei roto i te ripanga, ka huri haere tatou i roto i te huinga ka timata mai i te hash o te ki e rapu ana tatou. I ia mokowhiti, ka tirohia mena ko te matua te mea e rapu ana matou, a, ki te pera, ka whakahokia e matou tona uara. Ka tirohia ano e matou mehemea kei te putua te ki, a ki te pera, ka whakakorehia e matou te rapu.
Ki te kore e kitea te kī, ka whakahokia mai e te waehere he uara kore.
Ko enei mahinga rapu katoa ka taea te mahi i te wa kotahi ma te whakauru me te whakakore. Kei ia takirua o te ripanga tetahi o nga ahuatanga e wha kua whakaahuahia i runga ake mo te rere.
Te muku i roto i te ripanga hash
Waehere mo te whakakore i nga taviri:
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);
}
}
Ko te whakakore i tetahi taviri he mahi rerekee: ka waiho e tatou te ki ki te ripanga ka tohuhia tona uara (ehara i te ki tonu) he putua. He tino rite tenei waehere ki lookup(), engari ka kitea he taurite ki runga ki tetahi ki, ka noho kau tona uara.
Ka rite ki te korero i runga ake nei, ina tuhia he taviri ki te mokamoka, kare e nekehia. Ahakoa ka mukua tetahi huānga mai i te teepu, ka mau tonu te ki, ka noho kau noa tona uara. Ko te tikanga karekau e hiahia ki te whakamahi i te mahi tuhi ngota mo te uara mokamoka, na te mea kare he aha mena kei te putua te uara o naianei, kaore ranei - ka noho kau tonu.
Te rahi o te ripanga hash
Ka taea e koe te huri i te rahi o te ripanga hash ma te hanga i tetahi ripanga nui ake me te whakauru i nga huānga kore-kore mai i te ripanga tawhito ki roto. Kaore au i whakatinana i tenei mahi na te mea i pirangi ahau kia ngawari te tauira tauira. I tua atu, i roto i nga kaupapa CUDA, he maha nga wa ka mahia te tohatoha mahara i roto i te waehere kaihautu kaua i roto i te pata CUDA.
I roto i te tuhinga e whakaahua ana me pehea te whakarereke i taua hanganga raraunga kua parea-raka.
Te whakataetae
I roto i nga snippets waehere mahi i runga ake nei gpu_hashtable_insert(), _lookup() и _delete() Tukatia tetahi takirua uara-matua i ia wa. Na raro gpu_hashtable_insert(), _lookup() и _delete() tukatuka i te huinga takirua i roto i te whakarara, ia takirua ki tetahi miro mahi GPU motuhake:
// 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);
}
}
Ka tautoko te ripanga hash-aukati i nga whakaurunga, nga rapunga, me nga mukunga. Na te mea kei roto tonu nga takirua uara-matua i tetahi o nga kawanatanga e wha karekau nga ki e neke, ka whakamanahia e te ripanga te tika ahakoa ka whakamahia nga momo mahi i te wa kotahi.
Heoi, ki te tukatuka tatou i te puranga whakauru me te whakakore i te whakarara, a ki te mea he taviri taarua te huinga whakauru o nga takirua, kare e taea e tatou te matapae ko wai nga takirua ka "toa"—ka tuhia ki te ripanga hash whakamutunga. Me kii i karangahia e matou te waehere whakauru me te huinga whakaurunga takirua A/0 B/1 A/2 C/3 A/4. Ka oti te waehere, takirua B/1 и C/3 kua whakapumautia kei roto i te tepu, engari i te wa ano ka puta tetahi o nga takirua ki roto A/0, A/2 ranei A/4. He raruraru tenei, kaore ranei - kei runga i te tono. Kei te mohio koe i mua ake karekau he taviri taarua i te rarangi whakauru, kare pea koe e whakaaro ko tehea uara i tuhia whakamutunga.
Mena he raru tenei mo koe, me wehe e koe nga takirua takirua ki nga waea punaha CUDA rereke. I roto i te CUDA, ko nga mahi katoa e karanga ana i te pata ka oti tonu i mua i te waea kaara e whai ake nei (i te iti rawa i roto i te miro kotahi. I roto i nga miro rereke, ka mahia nga kakano ki te whakarara). I roto i te tauira i runga ake nei, ki te karanga koe i tetahi kernel me A/0 B/1 A/2 C/3, me tetahi atu me A/4, katahi te ki A ka whiwhi i te uara 4.
Inaianei me korero tatou mehemea me tika nga mahi lookup() и delete() whakamahi i te atatohu maamaa ranei ki te huinga takirua o te ripanga hash. E kii ana:
Ka taea e te kaikoipi te whiriwhiri ki te arotau i nga panui me te tuhi ki te maharatanga o te ao, te tiritahi ranei... Ka taea te whakakore i enei arotautanga ma te whakamahi i te kupu matua
volatile: ... ko nga tohutoro mo tenei taurangi ka whakahiatohia ki roto i te maharatanga tuturu mo te panui, tuhituhi ranei.
Karekau he tono mo nga whakaaro tika volatile. Mena ka whakamahia e te miro mahi he uara keteroki mai i te mahi panui o mua, katahi ka whakamahia nga korero kua tawhito. Heoi ano, he korero tenei mai i te ahua tika o te ripanga hash i tetahi wa o te karanga kernel. Mena ka hiahia koe ki te whakamahi i nga korero hou, ka taea e koe te whakamahi i te taurangi volatile, engari ka iti ake te mahi: e ai ki aku whakamatautau, i te wa i whakakorehia ai te 32 miriona huānga, ka heke te tere mai i te 500 miriona whakakore/hekona ki te 450 miriona whakakorenga/hekona.
Te whai hua
I roto i te whakamatautau mo te whakauru 64 miriona huānga me te whakakore i te 32 miriona o ratou, whakataetae i waenganui std::unordered_map a karekau he ripanga hash mo te GPU:
std::unordered_map i pau i te 70 ms ki te whakauru me te tango i nga huānga ka whakawāteahia unordered_map (Ko te whakakore i nga miriona o nga huānga he nui te wa, na te mea kei roto unordered_map he maha nga tohatoha mahara kua mahia). Te korero pono, std:unordered_map tino rereke nga here. He miro CPU kotahi mo te mahi, e tautoko ana i nga uara-matua ahakoa te rahi, he pai te mahi i runga i nga reiti whakamahi teitei, me te whakaatu i nga mahi pumau i muri i nga whakakorenga maha.
Ko te roanga o te ripanga hash mo te GPU me te whakawhitiwhiti-a-kaupapa ko te 984 ms. Kei roto i tenei ko te wa i whakapaua ki te whakatakoto i te teepu ki roto i te mahara me te muku (te tohatoha i te 1 GB o te mahara mo te wa kotahi, ka roa te wa i roto i te CUDA), te whakauru me te whakakore i nga huānga, me te huri i runga ake. Ko nga kape katoa mai i te maharatanga kaari ataata ka whai whakaaro hoki.
Ko te ripanga hash ake i 271 ms ki te whakaoti. Kei roto i tenei ko te wa e whakapaua ana e te kaari ataata ki te whakauru me te whakakore i nga huānga, karekau e whai whakaaro ki te wa i pau ki te kape ki te mahara me te huri i te ripanga ka puta. Mena kei te noho te tepu GPU mo te wa roa, mena kei te noho katoa te ripanga hash ki roto i te maharatanga o te kaari ataata (hei tauira, ki te hanga i tetahi ripanga hash ka whakamahia e etahi atu waehere GPU, kaua ko te tukatuka pokapū), katahi ka he tika te hua whakamatautau.
Ko te ripanga hash mo te kaari ataata e whakaatu ana i nga mahi teitei na te nui o te whakaputanga me te whakarara kaha.
hapehape
Ko te hoahoanga ripanga hash etahi take hei maarama:
- Ka whakararuraruhia te tirotiro raina e te rapoi, na te mea ka iti iho te whakatakoto o nga ki o te ripanga.
- Kaore e tangohia nga taviri ma te whakamahi i te mahi
deletea, i te roanga o te wa, ka marara te tepu.
Ko te mutunga mai, ka heke haere te mahinga o te ripanga hash, ina koa ka noho mo te wa roa, he maha nga whakaurunga me te whakakore. Ko tetahi huarahi ki te whakaiti i enei ngoikoretanga ko te whakahoki ano ki tetahi teepu hou me te iti o te reeti whakamahi me te tarai i nga kii kua tangohia i te wa o te rehashing.
Hei whakaatu i nga take kua whakaahuatia, ka whakamahi ahau i te waehere i runga ake nei ki te hanga tepu me te 128 miriona huānga me te huri i roto i te 4 miriona huānga kia whakakiia ra ano e ahau te 124 miriona mokowhiti (te 0,96 pea te utu whakamahi). Anei te ripanga hua, ko ia rarangi he piiraa CUDA kernel ki te whakauru i te 4 miriona huānga hou ki roto i te ripanga hash kotahi:
Te tere whakamahi
Te roanga whakauru 4 huānga
0,00
11,608448 ms (361,314798 miriona taviri/hekona)
0,03
11,751424 ms (356,918799 miriona taviri/hekona)
0,06
11,942592 ms (351,205515 miriona taviri/hekona)
0,09
12,081120 ms (347,178429 miriona taviri/hekona)
0,12
12,242560 ms (342,600233 miriona taviri/hekona)
0,16
12,396448 ms (338,347235 miriona taviri/hekona)
0,19
12,533024 ms (334,660176 miriona taviri/hekona)
0,22
12,703328 ms (330,173626 miriona taviri/hekona)
0,25
12,884512 ms (325,530693 miriona taviri/hekona)
0,28
13,033472 ms (321,810182 miriona taviri/hekona)
0,31
13,239296 ms (316,807174 miriona taviri/hekona)
0,34
13,392448 ms (313,184256 miriona taviri/hekona)
0,37
13,624000 ms (307,861434 miriona taviri/hekona)
0,41
13,875520 ms (302,280855 miriona taviri/hekona)
0,44
14,126528 ms (296,909756 miriona taviri/hekona)
0,47
14,399328 ms (291,284699 miriona taviri/hekona)
0,50
14,690304 ms (285,515123 miriona taviri/hekona)
0,53
15,039136 ms (278,892623 miriona taviri/hekona)
0,56
15,478656 ms (270,973402 miriona taviri/hekona)
0,59
15,985664 ms (262,379092 miriona taviri/hekona)
0,62
16,668673 ms (251,627968 miriona taviri/hekona)
0,66
17,587200 ms (238,486174 miriona taviri/hekona)
0,69
18,690048 ms (224,413765 miriona taviri/hekona)
0,72
20,278816 ms (206,831789 miriona taviri/hekona)
0,75
22,545408 ms (186,038058 miriona taviri/hekona)
0,78
26,053312 ms (160,989275 miriona taviri/hekona)
0,81
31,895008 ms (131,503463 miriona taviri/hekona)
0,84
42,103294 ms (99,619378 miriona taviri/hekona)
0,87
61,849056 ms (67,815164 miriona taviri/hekona)
0,90
105,695999 ms (39,682713 miriona taviri/hekona)
0,94
240,204636 ms (17,461378 miriona taviri/hekona)
Ka piki ake te whakamahinga, ka heke te mahi. I te nuinga o nga wa kaore tenei e pai. Mena ka whakaurua e tetahi tono nga huānga ki roto i te teepu katahi ka rukea (hei tauira, i te tatau kupu i roto i te pukapuka), ehara tenei i te raru. Engari ki te whakamahia e te tono he ripanga hash roa (hei tauira, i roto i te ētita whakairoiro ki te penapena i nga waahanga kore-kore o nga whakaahua ka whakauruhia e te kaiwhakamahi nga korero me te whakakore i nga korero), ka raru pea tenei whanonga.
Na ka ine i te hohonu o te riipene riipu i muri i te 64 miriona whakaurunga (tauwehe whakamahi 0,5). Ko te hohonutanga toharite ko te 0,4774, no reira ko te nuinga o nga taviri kei roto i te mokamoka pai rawa atu, kotahi te mokowhiti mai i te waahi pai rawa atu. Ko te hōhonutanga oro mōrahi he 60.
Ka ine ahau i te hohonutanga o te tirotiro i runga i te teepu me te 124 miriona whakaurunga (tauwehe whakamahi 0,97). Ko te hohonutanga toharite kua 10,1757, me te teitei - 6474 (!!). Ka tino heke te mahi taaro raina i nga reiti whakamahi teitei.
He pai ake kia noho iti te reiti whakamahi o tenei ripanga hash. Engari ka piki ake taatau mahi i runga i te whakapaunga o te kohi mahara. Waimarie, mo nga taviri moka-32 me nga uara, ka taea te whakatika. Mena i roto i te tauira i runga ake nei, i roto i te teepu me te 128 miriona nga huānga, ka pupuri tatou i te tauwehe whakamahi o te 0,25, katahi ka taea e tatou te whakanoho kia kaua e neke atu i te 32 miriona huānga ki roto, ka ngaro nga toenga 96 miriona mokamoka - 8 paita mo ia takirua. , 768 MB o te mahara ngaro.
Kia mahara kei te korero matou mo te ngaro o te mahara o te kaari ataata, he rauemi nui ake i te mahara punaha. Ahakoa ko te nuinga o nga kaari whakairoiro papamahi hou e tautoko ana i te CUDA he iti rawa te 4 GB o te mahara (i te wa e tuhi ana, ko te NVIDIA 2080 Ti he 11 GB), ehara tonu i te whakatau whakaaro nui kia ngaro aua moni.
A muri ake nei ka tuhia e au etahi atu korero mo te hanga ripanga hash mo nga kaari ataata kaore he raru mo te hohonu o te tirotiro, me nga huarahi ki te whakamahi ano i nga mokamoka kua mukua.
Te ine hohonutanga oro
Hei whakatau i te hōhonutanga tirotiro o te kī, ka taea e tatou te tango i te hash o te kī (tona taupū ripanga pai) mai i tana taupū ripanga tūturu:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Na te mahi makutu o nga tau takirua takirua a te tokorua me te mea ko te kaha o te ripanga hash e rua ki te kaha o te rua, ka whai hua tenei huarahi ahakoa ka nekehia te taurangi matua ki te timatanga o te ripanga. Me tango he taviri kua haua ki te 1, engari ka kuhu ki roto i te mokamoka 3. Na mo te teepu whai kaha 4 ka whiwhi tatou (3 — 1) & 3, he rite ki te 2.
mutunga
Mena he patai, he korero ranei, tena koa imeera mai ki ahau whakatuwhera kaupapa hou ranei ki roto.
I tuhia tenei waehere i raro i te faauruhia mai i nga tuhinga pai:
Hei nga ra kei te heke mai, ka tuhi tonu ahau mo nga whakatinanatanga ripanga hash mo nga kaari ataata me te tātari i o raatau mahi. Ko aku mahere ko te mekameka, te Robin Hood hashing, me te cuckoo hashing ma te whakamahi i nga mahi ngota i roto i nga hanganga raraunga e pai ana ki te GPU.
Source: will.com