Nilichapisha kwenye Github .
Ni jedwali rahisi la GPU lenye uwezo wa kuchakata mamia ya mamilioni ya viingilio kwa sekunde. Kwenye kompyuta yangu ndogo ya NVIDIA GTX 1060, msimbo huweka jozi milioni 64 za thamani-msingi zilizozalishwa bila mpangilio katika takriban ms 210 na kuondoa jozi milioni 32 kwa takriban 64 ms.
Hiyo ni, kasi kwenye kompyuta ndogo ni takriban milioni 300 kwa sekunde na kufuta milioni 500 kwa sekunde.
Jedwali limeandikwa katika CUDA, ingawa mbinu hiyo hiyo inaweza kutumika kwa HLSL au GLSL. Utekelezaji una vikwazo kadhaa ili kuhakikisha utendaji wa juu kwenye kadi ya video:
- Vifunguo vya 32-bit tu na maadili sawa yanachakatwa.
- Jedwali la hashi lina saizi isiyobadilika.
- Na ukubwa huu lazima uwe sawa na mbili kwa nguvu.
Kwa funguo na maadili, unahitaji kuhifadhi alama ya delimiter rahisi (katika nambari iliyo hapo juu hii ni 0xffffffff).
Jedwali la hash bila kufuli
Jedwali la hashi hutumia anwani wazi na , yaani, ni safu tu ya jozi za thamani-msingi ambazo zimehifadhiwa kwenye kumbukumbu na ina utendaji bora wa kache. Vile vile haviwezi kusemwa kwa minyororo, ambayo inahusisha kutafuta pointer katika orodha iliyounganishwa. Jedwali la hashi ni safu rahisi ya kuhifadhi vitu KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Saizi ya meza ni nguvu ya mbili, sio nambari kuu, kwa sababu maagizo moja ya haraka yanatosha kutumia pow2/AND mask, lakini mwendeshaji wa modulus ni polepole zaidi. Hii ni muhimu katika kesi ya uchunguzi wa mstari, kwa kuwa katika utafutaji wa jedwali la mstari index ya yanayopangwa lazima imefungwa katika kila slot. Na matokeo yake, gharama ya operesheni huongezwa modulo katika kila yanayopangwa.
Jedwali huhifadhi tu ufunguo na thamani kwa kila kipengele, sio heshi ya ufunguo. Kwa kuwa jedwali huhifadhi funguo 32-bit tu, heshi huhesabiwa haraka sana. Msimbo ulio hapo juu unatumia heshi ya Murmur3, ambayo hufanya mabadiliko machache tu, XOR na kuzidisha.
Jedwali la hashi hutumia mbinu za ulinzi za kufunga ambazo hazitegemei mpangilio wa kumbukumbu. Hata kama shughuli zingine za uandishi zitatatiza mpangilio wa shughuli zingine kama hizo, jedwali la hashi bado litadumisha hali sahihi. Tutazungumza juu ya hii hapa chini. Mbinu hiyo inafanya kazi vizuri na kadi za video zinazoendesha maelfu ya nyuzi kwa wakati mmoja.
Vifunguo na thamani katika jedwali la hashi huanzishwa hadi tupu.
Nambari inaweza kubadilishwa ili kushughulikia funguo na maadili ya 64-bit pia. Vifunguo vinahitaji shughuli za kusoma, kuandika na kulinganisha na kubadilishana kwa atomiki. Na maadili yanahitaji shughuli za kusoma na kuandika za atomiki. Kwa bahati nzuri, katika CUDA, shughuli za kusoma-kuandika kwa maadili ya 32- na 64-bit ni ya atomiki mradi tu zimepangwa asili (tazama hapa chini). ), na kadi za kisasa za video zinaunga mkono shughuli za 64-bit za atomiki za kulinganisha na kubadilishana. Bila shaka, wakati wa kuhamia bits 64, utendaji utapungua kidogo.
Hali ya jedwali la hash
Kila jozi ya thamani-msingi katika jedwali la heshi inaweza kuwa na mojawapo ya majimbo manne:
- Ufunguo na thamani ni tupu. Katika hali hii, jedwali la hashi limeanzishwa.
- Ufunguo umeandikwa, lakini thamani bado haijaandikwa. Ikiwa thread nyingine inasoma data kwa sasa, basi inarudi tupu. Hii ni ya kawaida, jambo lile lile lingetokea ikiwa nyuzi nyingine ya utekelezaji ingefanya kazi mapema kidogo, na tunazungumza juu ya muundo wa data unaofanana.
- Ufunguo na thamani zote mbili zimerekodiwa.
- Thamani inapatikana kwa nyuzi zingine za utekelezaji, lakini ufunguo bado. Hii inaweza kutokea kwa sababu muundo wa programu wa CUDA una mfano wa kumbukumbu ulioagizwa kwa urahisi. Hii ni kawaida; kwa hali yoyote, ufunguo bado hauna kitu, hata kama thamani haiko hivyo tena.
Nuance muhimu ni kwamba mara tu ufunguo umeandikwa kwenye slot, haisogei tena - hata ikiwa ufunguo umefutwa, tutazungumzia kuhusu hili hapa chini.
Nambari ya jedwali la hashi hata inafanya kazi na mifano ya kumbukumbu iliyopangwa kwa urahisi ambayo mpangilio ambao kumbukumbu inasomwa na kuandikwa haijulikani. Tunapoangalia uwekaji, utafutaji, na ufutaji katika jedwali la heshi, kumbuka kwamba kila jozi ya thamani-msingi iko katika mojawapo ya hali nne zilizoelezwa hapo juu.
Kuingiza kwenye jedwali la hashi
Kazi ya CUDA inayoingiza jozi za thamani-msingi kwenye jedwali la hashi inaonekana kama hii:
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);
}
}
Ili kuingiza kitufe, msimbo hurudia kupitia safu ya jedwali la heshi kuanzia na heshi ya kitufe kilichoingizwa. Kila nafasi katika safu hufanya operesheni ya kulinganisha na kubadilishana ya atomiki ambayo inalinganisha ufunguo katika nafasi hiyo na tupu. Ikigunduliwa kutolingana, ufunguo katika nafasi husasishwa kwa ufunguo ulioingizwa, na kisha ufunguo wa awali wa slot unarejeshwa. Ikiwa ufunguo huu wa asili ulikuwa tupu au ulilingana na ufunguo ulioingizwa, basi msimbo ulipata nafasi inayofaa kwa kuingizwa na kuingiza thamani iliyoingizwa kwenye slot.
Ikiwa katika simu moja ya kernel gpu_hashtable_insert() kuna vitu vingi vilivyo na ufunguo sawa, basi maadili yao yoyote yanaweza kuandikwa kwa slot muhimu. Hii inachukuliwa kuwa ya kawaida: moja ya thamani-ufunguo huandika wakati wa simu itafaulu, lakini kwa kuwa yote haya yanatokea kwa usawa ndani ya nyuzi kadhaa za utekelezaji, hatuwezi kutabiri ni kumbukumbu gani itakuwa ya mwisho.
Utaftaji wa jedwali la hash
Nambari ya funguo za kutafuta:
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);
}
}
Ili kupata thamani ya ufunguo uliohifadhiwa kwenye jedwali, tunarudia kupitia safu kuanzia na heshi ya kitufe tunachotafuta. Katika kila slot, tunaangalia ikiwa ufunguo ndio tunatafuta, na ikiwa ni hivyo, tunarudisha thamani yake. Pia tunaangalia ikiwa ufunguo hauna kitu, na ikiwa ni hivyo, tunabatilisha utafutaji.
Ikiwa hatuwezi kupata ufunguo, msimbo unarudi thamani tupu.
Shughuli hizi zote za utafutaji zinaweza kufanywa kwa wakati mmoja kwa kuingiza na kufuta. Kila jozi kwenye jedwali itakuwa na mojawapo ya majimbo manne yaliyoelezwa hapo juu kwa mtiririko.
Inafuta kwenye jedwali la hashi
Msimbo wa kufuta funguo:
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);
}
}
Kufuta ufunguo unafanywa kwa njia isiyo ya kawaida: tunaacha ufunguo kwenye meza na kuashiria thamani yake (sio ufunguo yenyewe) kuwa tupu. Kanuni hii inafanana sana na lookup(), isipokuwa kwamba mechi inapopatikana kwenye ufunguo, inafanya thamani yake kuwa tupu.
Kama ilivyoelezwa hapo juu, mara tu ufunguo umeandikwa kwenye slot, haisogezwi tena. Hata wakati kipengele kinafutwa kutoka kwa meza, ufunguo unabaki mahali, thamani yake inakuwa tupu. Hii ina maana kwamba hatuhitaji kutumia operesheni ya uandishi wa atomiki kwa thamani ya nafasi, kwa sababu haijalishi ikiwa thamani ya sasa ni tupu au la - bado itakuwa tupu.
Kubadilisha ukubwa wa jedwali la hashi
Unaweza kubadilisha ukubwa wa jedwali la hashi kwa kuunda jedwali kubwa zaidi na kuingiza vipengee visivyo tupu kutoka kwa jedwali kuu la zamani ndani yake. Sikutekeleza utendakazi huu kwa sababu nilitaka kuweka nambari ya sampuli kuwa rahisi. Zaidi ya hayo, katika programu za CUDA, ugawaji wa kumbukumbu mara nyingi hufanywa katika msimbo wa mwenyeji badala ya kernel ya CUDA.
Katika makala inaeleza jinsi ya kurekebisha muundo wa data uliolindwa kwa kufuli.
Ushindani
Katika vijisehemu vya msimbo wa kazi hapo juu gpu_hashtable_insert(), _lookup() ΠΈ _delete() chakata jozi moja ya thamani-msingi kwa wakati mmoja. Na chini gpu_hashtable_insert(), _lookup() ΠΈ _delete() chakata safu ya jozi sambamba, kila jozi kwenye uzi tofauti wa utekelezaji 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);
}
}
Jedwali la heshi linalostahimili kufuli linaauni uwekaji, utafutaji na ufutaji kwa wakati mmoja. Kwa sababu jozi za thamani-msingi huwa katika mojawapo ya majimbo manne na funguo hazisogei, jedwali huhakikisha usahihi hata wakati aina tofauti za utendakazi zinatumika kwa wakati mmoja.
Hata hivyo, tukichakata kundi la viambajengo na ufutaji kwa sambamba, na ikiwa safu ya ingizo ya jozi ina funguo rudufu, basi hatutaweza kutabiri ni jozi zipi "zitashinda" -itaandikwa kwenye jedwali la hashi mwishowe. Wacha tuseme tuliita msimbo wa kuingiza na safu ya pembejeo ya jozi A/0 B/1 A/2 C/3 A/4. Wakati msimbo ukamilika, jozi B/1 ΠΈ C/3 ni uhakika wa kuwepo katika meza, lakini wakati huo huo yoyote ya jozi itaonekana ndani yake A/0, A/2 au A/4. Hii inaweza kuwa shida au isiwe - yote inategemea programu. Unaweza kujua mapema kwamba hakuna funguo rudufu katika safu ya ingizo, au huenda usijali ni thamani gani iliyoandikwa mwisho.
Ikiwa hili ni tatizo kwako, basi unahitaji kutenganisha jozi mbili katika simu tofauti za mfumo wa CUDA. Katika CUDA, operesheni yoyote inayoita kernel daima hukamilika kabla ya simu inayofuata ya kernel (angalau ndani ya thread moja. Katika nyuzi tofauti, kernels hutekelezwa kwa usawa). Katika mfano hapo juu, ikiwa unaita kernel moja na A/0 B/1 A/2 C/3, na nyingine na A/4, kisha ufunguo A atapata thamani 4.
Sasa hebu tuzungumze juu ya ikiwa utendaji unapaswa lookup() ΠΈ delete() tumia kielekezi wazi au tete kwa safu ya jozi kwenye jedwali la hashi. Inasema kuwa:
Mkusanyaji anaweza kuchagua kuboresha usomaji na kuandika kwa kumbukumbu ya kimataifa au iliyoshirikiwa... Uboreshaji huu unaweza kulemazwa kwa kutumia neno kuu.
volatile: ... marejeleo yoyote ya utaftaji huu yanakusanywa katika kumbukumbu halisi ya kusoma au kuandika maagizo.
Mawazo ya usahihi hayahitaji maombi volatile. Ikiwa thread ya utekelezaji inatumia thamani iliyohifadhiwa kutoka kwa uendeshaji wa awali wa kusoma, basi itakuwa ikitumia maelezo ya zamani kidogo. Lakini bado, hii ni habari kutoka kwa hali sahihi ya jedwali la hashi wakati fulani wa simu ya kernel. Ikiwa unahitaji kutumia taarifa za hivi punde, unaweza kutumia faharasa volatile, lakini basi utendaji utapungua kidogo: kwa mujibu wa vipimo vyangu, wakati wa kufuta vipengele milioni 32, kasi ilipungua kutoka kwa kufuta milioni 500 / sec hadi milioni 450 kufuta / sec.
Uzalishaji
Katika mtihani wa kuingiza vipengele milioni 64 na kufuta milioni 32 kati yao, ushindani kati ya std::unordered_map na kwa kweli hakuna meza ya haraka ya GPU:
std::unordered_map ilitumia ms 70 kuingiza na kuondoa vipengee na kisha kuviweka huru unordered_map (kuondoa mamilioni ya vitu huchukua muda mwingi, kwa sababu ndani unordered_map ugawaji wa kumbukumbu nyingi hufanywa). Kusema ukweli, std:unordered_map vikwazo tofauti kabisa. Ni safu moja ya utekelezaji ya CPU, inasaidia maadili muhimu ya saizi yoyote, hufanya kazi vizuri kwa viwango vya juu vya utumiaji, na inaonyesha utendaji thabiti baada ya kufutwa mara nyingi.
Muda wa jedwali la heshi kwa GPU na mawasiliano kati ya programu ulikuwa 984 ms. Hii ni pamoja na muda uliotumika kuweka jedwali kwenye kumbukumbu na kuifuta (kutenga GB 1 ya kumbukumbu wakati mmoja, ambayo inachukua muda katika CUDA), kuingiza na kufuta vipengele, na kurudia juu yao. Nakala zote kwenda na kutoka kwa kumbukumbu ya kadi ya video pia huzingatiwa.
Jedwali la hashi lenyewe lilichukua ms 271 kukamilika. Hii inajumuisha muda uliotumiwa na kadi ya video kuingiza na kufuta vipengele, na haizingatii muda uliotumika kunakili kwenye kumbukumbu na kurudia juu ya jedwali linalosababisha. Ikiwa jedwali la GPU linaishi kwa muda mrefu, au ikiwa jedwali la hashi liko kabisa kwenye kumbukumbu ya kadi ya video (kwa mfano, kuunda jedwali la hashi ambalo litatumiwa na msimbo mwingine wa GPU na sio kichakataji cha kati), basi matokeo ya mtihani yanafaa.
Jedwali la hashi la kadi ya video linaonyesha utendakazi wa hali ya juu kwa sababu ya upitishaji wa juu na ulinganishaji unaofanya kazi.
Mapungufu
Usanifu wa jedwali la hashi una maswala machache ya kufahamu:
- Uchunguzi wa mstari unatatizwa na kuunganishwa, ambayo husababisha funguo kwenye jedwali kuwekwa chini ya ukamilifu.
- Funguo haziondolewa kwa kutumia chaguo la kukokotoa
deletena baada ya muda wao clutter meza.
Matokeo yake, utendaji wa jedwali la hashi unaweza kuharibika hatua kwa hatua, hasa ikiwa ipo kwa muda mrefu na ina kuingiza na kufuta nyingi. Njia moja ya kupunguza ubaya huu ni kurudisha nyuma jedwali jipya na kiwango cha chini cha matumizi na kuchuja vitufe vilivyoondolewa wakati wa kurejesha tena.
Ili kufafanua masuala yaliyoelezwa, nitatumia nambari iliyo hapo juu kuunda jedwali lenye vipengele milioni 128 na kupitisha vipengele milioni 4 hadi nitakapojaza nafasi milioni 124 (kiwango cha matumizi cha takriban 0,96). Hapa kuna jedwali la matokeo, kila safu ni simu ya CUDA ya kuingiza vitu vipya milioni 4 kwenye jedwali moja la hashi:
Kiwango cha matumizi
Muda wa uwekaji vipengele 4
0,00
11,608448 ms (funguo milioni 361,314798 kwa sekunde.)
0,03
11,751424 ms (funguo milioni 356,918799 kwa sekunde.)
0,06
11,942592 ms (funguo milioni 351,205515 kwa sekunde.)
0,09
12,081120 ms (funguo milioni 347,178429 kwa sekunde.)
0,12
12,242560 ms (funguo milioni 342,600233 kwa sekunde.)
0,16
12,396448 ms (funguo milioni 338,347235 kwa sekunde.)
0,19
12,533024 ms (funguo milioni 334,660176 kwa sekunde.)
0,22
12,703328 ms (funguo milioni 330,173626 kwa sekunde.)
0,25
12,884512 ms (funguo milioni 325,530693 kwa sekunde.)
0,28
13,033472 ms (funguo milioni 321,810182 kwa sekunde.)
0,31
13,239296 ms (funguo milioni 316,807174 kwa sekunde.)
0,34
13,392448 ms (funguo milioni 313,184256 kwa sekunde.)
0,37
13,624000 ms (funguo milioni 307,861434 kwa sekunde.)
0,41
13,875520 ms (funguo milioni 302,280855 kwa sekunde.)
0,44
14,126528 ms (funguo milioni 296,909756 kwa sekunde.)
0,47
14,399328 ms (funguo milioni 291,284699 kwa sekunde.)
0,50
14,690304 ms (funguo milioni 285,515123 kwa sekunde.)
0,53
15,039136 ms (funguo milioni 278,892623 kwa sekunde.)
0,56
15,478656 ms (funguo milioni 270,973402 kwa sekunde.)
0,59
15,985664 ms (funguo milioni 262,379092 kwa sekunde.)
0,62
16,668673 ms (funguo milioni 251,627968 kwa sekunde.)
0,66
17,587200 ms (funguo milioni 238,486174 kwa sekunde.)
0,69
18,690048 ms (funguo milioni 224,413765 kwa sekunde.)
0,72
20,278816 ms (funguo milioni 206,831789 kwa sekunde.)
0,75
22,545408 ms (funguo milioni 186,038058 kwa sekunde.)
0,78
26,053312 ms (funguo milioni 160,989275 kwa sekunde.)
0,81
31,895008 ms (funguo milioni 131,503463 kwa sekunde.)
0,84
42,103294 ms (funguo milioni 99,619378 kwa sekunde.)
0,87
61,849056 ms (funguo milioni 67,815164 kwa sekunde.)
0,90
105,695999 ms (funguo milioni 39,682713 kwa sekunde.)
0,94
240,204636 ms (funguo milioni 17,461378 kwa sekunde.)
Kadiri matumizi yanavyoongezeka, utendaji hupungua. Hii haipendi katika hali nyingi. Ikiwa programu itaingiza vipengele kwenye meza na kisha kuvitupa (kwa mfano, wakati wa kuhesabu maneno katika kitabu), basi hii sio tatizo. Lakini ikiwa programu hutumia jedwali la hashi la muda mrefu (kwa mfano, katika kihariri cha picha kuhifadhi sehemu zisizo tupu za picha ambapo mtumiaji mara nyingi huingiza na kufuta habari), basi tabia hii inaweza kuwa tatizo.
Na kupima kina cha uchunguzi wa jedwali la heshi baada ya kuingiza milioni 64 (kipengele cha matumizi 0,5). Kina cha wastani kilikuwa 0,4774, kwa hivyo funguo nyingi zilikuwa kwenye nafasi nzuri zaidi au nafasi moja mbali na nafasi bora zaidi. Kina cha juu cha sauti kilikuwa 60.
Kisha nikapima kina cha uchunguzi kwenye jedwali na viingilio milioni 124 (sababu ya utumiaji 0,97). Kina cha wastani kilikuwa tayari 10,1757, na kiwango cha juu - 6474 (!!). Utendaji wa hisia za mstari hushuka sana kwa viwango vya juu vya utumiaji.
Ni bora kupunguza kiwango cha matumizi ya jedwali hili la hashi. Lakini basi tunaongeza utendaji kwa gharama ya matumizi ya kumbukumbu. Kwa bahati nzuri, katika kesi ya funguo 32-bit na maadili, hii inaweza kuhesabiwa haki. Ikiwa katika mfano hapo juu, kwenye jedwali iliyo na vitu milioni 128, tunaweka kipengele cha utumiaji cha 0,25, basi hatuwezi kuweka vitu zaidi ya milioni 32 ndani yake, na nafasi zilizobaki milioni 96 zitapotea - ka 8 kwa kila jozi. , 768 MB ya kumbukumbu iliyopotea.
Tafadhali kumbuka kuwa tunazungumzia juu ya kupoteza kumbukumbu ya kadi ya video, ambayo ni rasilimali ya thamani zaidi kuliko kumbukumbu ya mfumo. Ingawa kadi nyingi za kisasa za picha za mezani zinazounga mkono CUDA zina kumbukumbu ya angalau 4 GB (wakati wa kuandika, NVIDIA 2080 Ti ina GB 11), bado haingekuwa uamuzi wa busara zaidi kupoteza kiasi kama hicho.
Baadaye nitaandika zaidi juu ya kuunda meza za hashi kwa kadi za video ambazo hazina shida na kina cha uchunguzi, na pia njia za kutumia tena nafasi zilizofutwa.
Kipimo cha kina cha sauti
Kuamua kina cha uchunguzi wa ufunguo, tunaweza kutoa heshi ya ufunguo (faharisi yake bora ya jedwali) kutoka kwa faharisi yake halisi ya jedwali:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Kwa sababu ya uchawi wa nambari za binary zinazosaidia mbili na ukweli kwamba uwezo wa jedwali la hashi ni mbili kwa nguvu ya mbili, mbinu hii itafanya kazi hata wakati index muhimu inapohamishwa hadi mwanzo wa meza. Wacha tuchukue ufunguo ulioharakishwa hadi 1, lakini umeingizwa kwenye slot 3. Kisha kwa meza yenye uwezo wa 4 tunapata. (3 β 1) & 3, ambayo ni sawa na 2.
Hitimisho
Ikiwa una maswali au maoni, tafadhali nitumie barua pepe kwa au fungua mada mpya ndani.
Nambari hii iliandikwa chini ya msukumo kutoka kwa nakala bora:
Katika siku zijazo, nitaendelea kuandika kuhusu utekelezaji wa meza ya hash kwa kadi za video na kuchambua utendaji wao. Mipango yangu ni pamoja na minyororo, Robin Hood hashing, na cuckoo hashing kutumia shughuli za atomiki katika miundo ya data ambayo ni rafiki wa GPU.
Chanzo: mapenzi.com