Ndakaiisa paGithub chirongwa chitsva A Nyore GPU Hash Tafura.

Iyo iri nyore GPU hashi tafura inokwanisa kugadzirisa mazana emamiriyoni ekuisa pasekondi. Pane yangu NVIDIA GTX 1060 laptop, iyo kodhi inoisa 64 miriyoni yakangogadzirwa kiyi-kukosha mapeya mune angangoita 210 ms uye inobvisa 32 miriyoni pairi mune angangoita makumi matanhatu nenhatu ms.

Kureva kuti, kumhanya palaptop kunosvika mazana matatu emamiriyoni ekuisa/sec uye 300 miriyoni yekubvisa/sec.

Tafura yakanyorwa muCUDA, kunyangwe nzira imwechete inogona kushandiswa kuHLSL kana GLSL. Kuitwa kwacho kune zvipingamupinyi zvakati kuti ive nechokwadi chekuita kwepamusoro pane kadhi revhidhiyo:

Chete 32-bit makiyi uye maitiro akafanana anogadziriswa.
Tafura yehashi ine saizi yakagadziriswa.
Uye saizi iyi inofanira kuenzana nembiri kune simba.

Kune makiyi uye kukosha, iwe unofanirwa kuchengetedza yakapusa delimiter marker (mune kodhi iri pamusoro iyi 0xffffffff).

Hash tafura isina makiyi

Iyo hashi tafura inoshandisa yakavhurika kero ne linear probing, ndiko kuti, ingori rondedzero yekiyi-value pairs inochengetwa mundangariro uye ine yepamusoro cache performance. Izvo zvakafanana hazvigone kutaurwa nezvechaining, iyo inosanganisira kutsvaga chinongedzo mune yakabatana runyorwa. Tafura yehashi iri nyore hurongwa hwekuchengeta zvinhu KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Saizi yetafura isimba remaviri, kwete nhamba yepamusoro, nekuti kuraira kumwe kwekukurumidza kwakakwana kuisa iyo pow2/AND mask, asi modulus opareta inononoka. Izvi zvakakosha munyaya yekuongororwa kwemutsara, sezvo mumutsara wetafura yekutarisa iyo slot index inofanirwa kuputirwa mune yega yega slot. Uye semhedzisiro, mutengo wekushanda unowedzerwa modulo mune yega yega slot.

Tafura inongochengeta kiyi uye kukosha kwechimwe nechimwe chinhu, kwete hashi yekiyi. Sezvo tafura ichingochengeta 32-bit kiyi, iyo hashi inoverengerwa nekukurumidza. Iyo kodhi iri pamusoro inoshandisa iyo Murmur3 hashi, iyo inongoita mashoma mashifiti, XOR uye kuwanda.

Iyo tafura yehashi inoshandisa nzira dzekudzivirira dzekuvhara iyo yakazvimirira yehurongwa hwekurangarira. Kunyangwe mamwe mabasa ekunyora achikanganisa kurongeka kwemamwe maitirwo akadai, tafura yehashi icharamba ichichengetedza mamiriro chaiwo. Tichataura pamusoro peizvi pasi apa. Iyo tekinoroji inoshanda zvikuru nemakadhi evhidhiyo anomhanyisa zviuru zveshinda panguva imwe chete.

Makiyi uye kukosha muhashi tafura inotangwa kuti ishaye.

Iyo kodhi inogona kugadziridzwa kubata 64-bit makiyi uye kukosha zvakare. Makiyi anoda kuverenga, kunyora, uye kuenzanisa-uye-kuchinjanisa maatomu. Uye kukosha kunoda atomic kuverenga nekunyora mashandiro. Sezvineiwo, muCUDA, kuverenga-kunyora mashandiro e32- uye 64-bit maitiro ari atomu chero achingoenderana (ona pazasi). pano), uye makadhi emazuva ano evhidhiyo anotsigira 64-bit atomic kuenzanisa-uye-kuchinjana mashandiro. Ehe, kana uchienda kune 64 bits, kuita kunoderera zvishoma.

Hash table state

Imwe neimwe kiyi-kukosha pairi muhashi tafura inogona kuve neimwe yezvina nyika:

Kiyi uye kukosha hazvina chinhu. Mune ino mamiriro, iyo hashi tafura inotangwa.
Kiyi yakanyorwa pasi, asi kukosha hakusati kwanyorwa. Kana imwe tambo iri kuverenga data parizvino, inobva yadzoka isina chinhu. Izvi zvakajairika, chinhu chimwe chete chingadai chakaitika dai imwe tambo yekuuraya yakashanda zvishoma kare, uye isu tiri kutaura nezve yakabatana data chimiro.
Zvose kiyi uye kukosha zvinonyorwa.
Iko kukosha kunowanikwa kune dzimwe tambo dzekuuraya, asi kiyi haisati yasvika. Izvi zvinogona kuitika nekuti iyo CUDA programming modhi ine yakasununguka yakarairwa ndangariro modhi. Izvi zvakajairika; chero chiitiko, kiyi ichiri isina chinhu, kunyangwe kukosha kwacho kusisiri kudaro.

Chinhu chakakosha nuance ndechekuti kana kiyi yanyorwa kune slot, haichafambi - kunyangwe kiyi ikabviswa, isu tichataura pamusoro peizvi pazasi.

Iyo hash tafura kodhi inotoshanda neakasununguka akaodha memodheru mamodheru umo marongero ayo ndangariro inoverengwa nekunyorwa isingazivikanwe. Sezvatinotarisa kuisa, kutarisa, uye kudzima mutafura yehashi, rangarira kuti imwe neimwe kiyi-kukosha peya iri mune imwe yematunhu mana atsanangurwa pamusoro.

Kupinza muhashi tafura

Basa reCUDA rinoisa kiyi-kukosha pairi mutafura yehashi rinotaridzika seizvi:

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

Kuisa kiyi, iyo kodhi inodzokorora kuburikidza nehashi tafura array kutanga nehashi yekiyi yakaiswa. Imwe neimwe slot muhurongwa inoita atomic kuenzanisa-uye-kuchinjanisa oparesheni inofananidza kiyi iri muslot iyoyo kuti isina. Kana kusapindirana kwaonekwa, kiyi mu slot inovandudzwa nekiyi yakaiswa, uye ipapo kiyi yekutanga slot inodzoserwa. Kana kiyi iyi yepakutanga yakanga isina kana kuenderana nekiyi yakaiswa, ipapo kodhi yakawana slot yakakodzera yekuisa uye yakaisa kukosha kwakaiswa mukati meiyo slot.

Kana mune imwe kernel kufona gpu_hashtable_insert() kune akati wandei zvinhu zvine kiyi imwechete, saka chero ipi zvayo yakakosha inogona kunyorwa kune kiyi slot. Izvi zvinoonekwa sezvakajairwa: imwe yekiyi-kukosha inonyora panguva yekufona ichabudirira, asi sezvo zvese izvi zvichiitika zvakafanana mukati mematanho akati wandei ekuuraya, isu hatigone kufanotaura kuti ndeipi ndangariro inonyora ichave yekupedzisira.

Hash tafura yekutarisa

Kodhi yekutsvaga 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 tiwane kukosha kwekiyi yakachengetwa mutafura, tinodzokorora kuburikidza nehurongwa kutanga nehashi yekiyi yatiri kutsvaga. Muchikamu chega chega, tinotarisa kana kiyi ndiyo yatiri kutsvaga, uye kana zvakadaro, tinodzorera kukosha kwayo. Isu tinoongorora zvakare kana kiyi isina chinhu, uye kana zvakadaro, tinobvisa kutsvaga.

Kana isu tisingakwanisi kuwana kiyi, iyo kodhi inodzosa kukosha kusina chinhu.

Ese aya mabasa ekutsvaga anogona kuitwa panguva imwe chete kuburikidza nekuisa uye kudzima. Peya yega yega patafura ichange iine imwe yematunhu mana anotsanangurwa pamusoro pekuyerera.

Kudzima mutafura yehashi

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

Kubvisa kiyi kunoitwa nenzira isina kujairika: tinosiya kiyi mutafura uye tinomaka kukosha kwayo (kwete kiyi pachayo) sechinhu chisina chinhu. Iyi kodhi yakafanana ne lookup(), kunze kwekuti kana machisi awanikwa pakiyi, anoita kuti kukosha kwawo kuve kusina.

Sezvataurwa pamusoro apa, kana kiyi yanyorwa kune slot, haichafambiswi. Kunyangwe kana chinhu chadzimwa patafura, kiyi inoramba iripo, kukosha kwayo kunongove isina chinhu. Izvi zvinoreva kuti hatifanire kushandisa atomic kunyora mashandiro kune iyo slot kukosha, nekuti hazvina basa kuti kukosha kwazvino hakuna chinhu kana kwete - icharamba isina chinhu.

Kugadzirisa tafura yehashi

Iwe unogona kushandura saizi yetafura yehashi nekugadzira tafura yakakura uye nekuisa zvinhu zvisina chinhu kubva patafura yekare mairi. Ini handina kuita basa iri nekuti ndaida kuchengeta kodhi yemuenzaniso iri nyore. Uyezve, muzvirongwa zveCUDA, kugoverwa kwendangariro kunowanzoitwa mukodhi yekutambira kwete muCUDA kernel.

Chinyorwa A Lock-Yemahara Wait-Yemahara Hash Tafura inotsanangura maitiro ekugadzirisa akadaro evhavha-yakachengetedzwa data chimiro.

Kukwikwidzana

Mune iri pamusoro pebasa rekodhi snippets gpu_hashtable_insert(), _lookup() и _delete() gadzira imwe kiyi-value peya panguva. Uye pasi gpu_hashtable_insert(), _lookup() и _delete() gadzira rondedzero yezviviri zvakafanana, peya yega yega mune yakaparadzana GPU tambo yekuuraya:

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

Iyo yekuvhara-inodzivirira hashi tafura inotsigira pamwe chete kuisa, kutarisa, uye kudzima. Nekuti makiyi-kukosha mapeya anogara ari mune imwe yematunhu mana uye makiyi haafambi, tafura inovimbisa kurongeka kunyangwe mhando dzakasiyana dzekushanda dzichishandiswa panguva imwe chete.

Zvisineyi, kana tikagadzirisa bhechi rekuisa nekudzimurwa nenzira yakafanana, uye kana rondedzero yezviviri iine makiyi akadzokororwa, hatizokwanisi kufanotaura kuti ndeapi maviri maviri acha “hwina”—achazonyorerwa kutafura yehashi kekupedzisira. Ngatitii takadaidza kodhi yekuisa nerunyoro rwemapeya A/0 B/1 A/2 C/3 A/4. Kana kodhi yapera, paviri B/1 и C/3 vanovimbiswa kuvapo mutafura, asi panguva imwechete chero ipi zvayo yevaviri ichaonekwa mairi A/0, A/2 kana A/4. Izvi zvinogona kana kusave dambudziko - zvese zvinoenderana nekushandisa. Unogona kuziva pachine nguva kuti hapana makiyi akadhindwa muchikamu chekuisa, kana kuti ungasave nehanya kuti ndeupi kukosha kwakanyorwa pekupedzisira.

Kana iri riri dambudziko kwauri, saka iwe unofanirwa kupatsanura iwo maviri maviri maviri muakasiyana CUDA system mafoni. MuCUDA, chero kuvhiya kunodaidza kernel nguva dzose kunopera kusati kwafona kernel inotevera (pamwe mukati meshinda imwe chete. Mushinda dzakasiyana, kernels dzinoitwa dzakafanana). Mumuenzaniso uri pamusoro, kana iwe ukadaidza imwe kernel nayo A/0 B/1 A/2 C/3, uye mumwe wacho A/4, ipapo kiyi A uchawana kukosha 4.

Zvino ngatitaure pamusoro pekuti mabasa anofanira lookup() и delete() shandisa plain or volatile pointer kune akatevedzana peya patafura yehashi. CUDA Zvinyorwa Inoti:

Muumbi anogona kusarudza kukwenenzvera kuverenga nekunyora kundangariro dzepasi rose kana dzakagovaniswa... Izvi zvigadziriso zvinogona kudzimwa uchishandisa izwi rakakosha. volatile: ... chero chirevo chekuchinja uku chinounganidzwa mundangariro chaiyo kuverenga kana kunyora rairo

Kururamisa kufunga hakudi kushandiswa volatile. Kana iyo tambo yekuuraya ichishandisa cached kukosha kubva kune yakare kuverenga oparesheni, saka inenge ichishandisa ruzivo rwechinyakare. Asi zvakadaro, urwu ruzivo kubva kune chaiyo mamiriro etafura yehashi pane imwe nguva yekufona kernel. Kana uchida kushandisa ruzivo rwazvino, unogona kushandisa index volatile, asi ipapo kuita kuchadzikira zvishoma: zvinoenderana nemiedzo yangu, pakudzima 32 miriyoni zvinhu, kumhanya kwakadzikira kubva pa500 miriyoni deletions/sec kusvika 450 miriyoni deletions/sec.

Kubudirira

Muyedzo yekuisa 64 miriyoni zvinhu uye kudzima 32 miriyoni yacho, makwikwi pakati std::unordered_map uye hapana kana hashi tafura yeGPU:

std::unordered_map akashandisa 70 ms kuisa nekubvisa zvinhu uye nekuzvisunungura unordered_map (kubvisa mamirioni ezvimiro zvinotora nguva yakawanda, nekuti mukati unordered_map kugoverwa kwendangariro kwakawanda kunoitwa). Kutaura chokwadi, std:unordered_map zvirambidzo zvakasiyana zvachose. Iyo imwechete CPU tambo yekuuraya, inotsigira kiyi-yakakosha chero saizi, inoita zvakanaka pamazinga ekushandisa akakwira, uye inoratidza kuita kwakagadzikana mushure mekudzima kwakawanda.

Nguva yetafura yehashi yeGPU uye yepakati-chirongwa kutaurirana yaive 984 ms. Izvi zvinosanganisira nguva yakashandiswa kuisa tafura mundangariro uye kuidzima (kugovera 1 GB yekuyeuka imwe nguva, izvo zvinotora nguva muCUDA), kuisa nekudzima zvinhu, uye kudzokorodza pamusoro pazvo. Ese makopi kuenda uye kubva kune vhidhiyo kadhi memory anotariswa zvakare.

Iyo hashi tafura pachayo yakatora 271 ms kuti ipedze. Izvi zvinosanganisira nguva inopedzwa nekadhi revhidhiyo kuisa nekudzima zvinhu, uye hazvifungidzire nguva yakashandiswa kukopa mundangariro nekudzokorora pamusoro petafura inoguma. Kana iyo tafura yeGPU ichirarama kwenguva yakareba, kana kana tafura yehashi irimo zvachose mundangariro yevhidhiyo kadhi (semuenzaniso, kugadzira tafura yehashi iyo ichashandiswa neimwe GPU kodhi uye kwete yepakati processor), ipapo mhinduro yebvunzo yakakosha.

Tafura yehashi yekadhi yevhidhiyo inoratidza kushanda kwepamusoro nekuda kwekuwedzera kwepamusoro uye kushanda kwekuenzanisa.

kutadza

Iyo hash tafura yekuvaka ine mashoma mashoma ekuziva:

Linear probing inokanganiswa nekubatanidza, izvo zvinoita kuti makiyi ari mutafura aiswe zvishoma pane zvakakwana.
Makiyi haana kubviswa uchishandisa basa delete uye nekufamba kwenguva vanokanganisa tafura.

Nekuda kweizvozvo, kuita kwetafura yehashi kunogona kuderera zvishoma nezvishoma, kunyanya kana iripo kwenguva yakareba uye iine akawanda anoisa uye anodzima. Imwe nzira yekudzikisira zvimhingamupinyi izvi ndeyekudzosera mutafura nyowani ine mwero wakaderera wekushandisa uye kusefa makiyi akabviswa panguva yekudzokorora.

Kuenzanisira nyaya dzakatsanangurwa, ini ndichashandisa kodhi iri pamusoro kugadzira tafura ine 128 miriyoni zvinhu uye loop kuburikidza nemamiriyoni mana zvinhu kusvika ndazadza 4 miriyoni slots (yekushandisa chiyero ingangoita 124). Heino tafura yemhedzisiro, mutsara wega wega ndeye CUDA kernel kufona kuti uise mamirioni mana zvinhu zvitsva mune imwe hashi tafura:

Mutengo wekushandisa
Nguva yekuiswa 4 zvinhu

0,00
11,608448 ms (361,314798 miriyoni makiyi/sec.)

0,03
11,751424 ms (356,918799 miriyoni makiyi/sec.)

0,06
11,942592 ms (351,205515 miriyoni makiyi/sec.)

0,09
12,081120 ms (347,178429 miriyoni makiyi/sec.)

0,12
12,242560 ms (342,600233 miriyoni makiyi/sec.)

0,16
12,396448 ms (338,347235 miriyoni makiyi/sec.)

0,19
12,533024 ms (334,660176 miriyoni makiyi/sec.)

0,22
12,703328 ms (330,173626 miriyoni makiyi/sec.)

0,25
12,884512 ms (325,530693 miriyoni makiyi/sec.)

0,28
13,033472 ms (321,810182 miriyoni makiyi/sec.)

0,31
13,239296 ms (316,807174 miriyoni makiyi/sec.)

0,34
13,392448 ms (313,184256 miriyoni makiyi/sec.)

0,37
13,624000 ms (307,861434 miriyoni makiyi/sec.)

0,41
13,875520 ms (302,280855 miriyoni makiyi/sec.)

0,44
14,126528 ms (296,909756 miriyoni makiyi/sec.)

0,47
14,399328 ms (291,284699 miriyoni makiyi/sec.)

0,50
14,690304 ms (285,515123 miriyoni makiyi/sec.)

0,53
15,039136 ms (278,892623 miriyoni makiyi/sec.)

0,56
15,478656 ms (270,973402 miriyoni makiyi/sec.)

0,59
15,985664 ms (262,379092 miriyoni makiyi/sec.)

0,62
16,668673 ms (251,627968 miriyoni makiyi/sec.)

0,66
17,587200 ms (238,486174 miriyoni makiyi/sec.)

0,69
18,690048 ms (224,413765 miriyoni makiyi/sec.)

0,72
20,278816 ms (206,831789 miriyoni makiyi/sec.)

0,75
22,545408 ms (186,038058 miriyoni makiyi/sec.)

0,78
26,053312 ms (160,989275 miriyoni makiyi/sec.)

0,81
31,895008 ms (131,503463 miriyoni makiyi/sec.)

0,84
42,103294 ms (99,619378 miriyoni makiyi/sec.)

0,87
61,849056 ms (67,815164 miriyoni makiyi/sec.)

0,90
105,695999 ms (39,682713 miriyoni makiyi/sec.)

0,94
240,204636 ms (17,461378 miriyoni makiyi/sec.)

Sezvo kushandiswa kunowedzera, kushanda kunoderera. Izvi hazvidiwi kazhinji. Kana application ikaisa zvinhu mutafura yobva yazvirasa (somuenzaniso, pakuverenga mazwi mubhuku), iri harisi dambudziko. Asi kana iyo application ikashandisa tafura yenguva refu yehashi (semuenzaniso, mune yemifananidzo mupepeti kuchengetedza zvikamu zvisina chinhu zvemifananidzo apo mushandisi anogara achiisa nekudzima ruzivo), ipapo maitiro aya anogona kunetsa.

Uye yakayera iyo hashi tafura yekuongorora kudzika mushure meiyo 64 miriyoni kuisa (kushandiswa chinhu 0,5). Hudzamu hwepakati hwaive 0,4774, saka makiyi mazhinji aive mune yakanakisa inogoneka slot kana imwe slot kure nenzvimbo yakanakisa. Kudzika kwekurira kwakanyanya kwaive makumi matanhatu.

Ini ndakabva ndayera kudzika kwekuongorora patafura ine 124 miriyoni yekuisa (kushandisa chinhu 0,97). Avhareji yakadzika yaive yatove 10,1757, uye iyo yakanyanya - 6474 (!!). Linear sensing performance inodonha zvakanyanya pamazinga ekushandisa akakwira.

Zvakanakisa kuchengetedza iyi hashi tafura yekushandisa yakaderera. Asi ipapo isu tinowedzera mashandiro pamutengo wekushandisa ndangariro. Sezvineiwo, kana iri 32-bit kiyi uye kukosha, izvi zvinogona kururamiswa. Kana mumuenzaniso wepamusoro, mutafura ine 128 miriyoni zvinhu, isu takachengeta utilization factor ye0,25, saka isu hatigone kuisa zvisingapfuure 32 miriyoni zvinhu mairi, uye yasara 96 miriyoni slots icharasika - 8 bytes pairi imwe neimwe. , 768 MB yekurasika ndangariro.

Ndapota cherechedza kuti isu tiri kutaura pamusoro pekurasikirwa kwevhidhiyo kadhi memory, iyo inonyanya kukosha sosi kupfuura system memory. Kunyangwe mazhinji emazuva ano desktop graphics makadhi anotsigira CUDA ane angangoita 4 GB yendangariro (panguva yekunyora, iyo NVIDIA 2080 Ti ine 11 GB), haingave isati iri iyo yakangwara sarudzo yekurasikirwa nemari yakadaro.

Gare gare ini ndichanyora zvimwe nezve kugadzira hashi matafura emakadhi evhidhiyo asina matambudziko nekudzika kwekuongorora, pamwe nenzira dzekushandisa zvakare dzakadzimwa slots.

Kurira kudzika kuyerwa

Kuti uone kudzika kwekiyi, tinogona kubvisa hashi yekiyi (yakakodzera tafura index) kubva kune yayo chaiyo tafura index:

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

Nekuda kwemashiripiti evaviri maviri anoenderana nhamba dzebhinari uye chokwadi chekuti kugona kwetafura yehashi kune maviri kune simba rezviviri, nzira iyi ichashanda kunyangwe iyo kiyi index inotamiswa kusvika pakutanga kwetafura. Ngatitorei kiyi ine hashi kuna 1, asi inoiswa mu slot 3. Zvino yetafura ine huwandu 4 tinowana (3 — 1) & 3, iyo yakaenzana ne2.

mhedziso

Kana uine mibvunzo kana mhinduro, ndapota nditumire email pa Twitter kana kuvhura dingindira idzva mukati repositories.

Iyi kodhi yakanyorwa pasi pekurudziro kubva kuzvinyorwa zvakanakisa:

Mune ramangwana, ini ndicharamba ndichinyora nezve hash tafura yekushandiswa kwemakadhi evhidhiyo uye kuongorora maitiro avo. Zvirongwa zvangu zvinosanganisira chaining, Robin Hood hashing, uye cuckoo hashing uchishandisa maatomu mashandiro mune data zvimiro zvine hushamwari GPU.

Source: www.habr.com

Rakareruka hash tafura yeGPU