Na-post ko ito sa Github .
Ito ay isang simpleng GPU hash table na may kakayahang magproseso ng daan-daang milyong mga pagsingit bawat segundo. Sa aking NVIDIA GTX 1060 laptop, ang code ay naglalagay ng 64 milyon na random na nabuong key-value pairs sa humigit-kumulang 210 ms at nag-aalis ng 32 milyong pares sa halos 64 ms.
Ibig sabihin, ang bilis sa isang laptop ay humigit-kumulang 300 million inserts/sec at 500 million deletes/sec.
Ang talahanayan ay nakasulat sa CUDA, bagaman ang parehong pamamaraan ay maaaring ilapat sa HLSL o GLSL. Ang pagpapatupad ay may ilang mga limitasyon upang matiyak ang mataas na pagganap sa isang video card:
- Ang mga 32-bit na key lamang at ang parehong mga halaga ang naproseso.
- Ang hash table ay may nakapirming laki.
- At ang sukat na ito ay dapat na katumbas ng dalawa sa kapangyarihan.
Para sa mga key at value, kailangan mong magreserba ng isang simpleng delimiter marker (sa code sa itaas ito ay 0xffffffff).
Hash table na walang lock
Ang hash table ay gumagamit ng open addressing with , ibig sabihin, isa lamang itong hanay ng mga pares ng key-value na nakaimbak sa memorya at may mahusay na pagganap ng cache. Ang parehong ay hindi maaaring sabihin para sa chaining, na nagsasangkot ng paghahanap para sa isang pointer sa isang naka-link na listahan. Ang hash table ay isang simpleng array na nag-iimbak ng mga elemento KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Ang laki ng talahanayan ay isang kapangyarihan ng dalawa, hindi isang pangunahing numero, dahil ang isang mabilis na pagtuturo ay sapat na upang ilapat ang pow2/AND mask, ngunit ang modulus operator ay mas mabagal. Ito ay mahalaga sa kaso ng linear probing, dahil sa isang linear table lookup ang slot index ay dapat na nakabalot sa bawat slot. At bilang isang resulta, ang gastos ng operasyon ay idinagdag modulo sa bawat slot.
Ang talahanayan ay nag-iimbak lamang ng susi at halaga para sa bawat elemento, hindi isang hash ng susi. Dahil ang talahanayan ay nag-iimbak lamang ng mga 32-bit na key, ang hash ay kinakalkula nang napakabilis. Ang code sa itaas ay gumagamit ng Murmur3 hash, na gumaganap lamang ng ilang shift, XOR at multiplikasyon.
Gumagamit ang hash table ng mga diskarte sa proteksyon sa pag-lock na hindi nakasalalay sa pagkakasunud-sunod ng memorya. Kahit na ang ilang mga write operations ay nakakagambala sa pagkakasunud-sunod ng iba pang mga naturang operasyon, ang hash table ay mananatili pa rin sa tamang estado. Pag-uusapan natin ito sa ibaba. Ang pamamaraan ay mahusay na gumagana sa mga video card na nagpapatakbo ng libu-libong mga thread nang sabay-sabay.
Ang mga susi at halaga sa hash table ay sinisimulan sa walang laman.
Maaaring mabago ang code upang mahawakan din ang mga 64-bit na key at value. Ang mga key ay nangangailangan ng atomic read, write, at compare-and-swap operations. At ang mga halaga ay nangangailangan ng atomic read and write operations. Sa kabutihang palad, sa CUDA, ang mga pagpapatakbo ng read-write para sa 32- at 64-bit na mga halaga ay atomic hangga't natural na nakahanay ang mga ito (tingnan sa ibaba). ), at ang mga modernong video card ay sumusuporta sa 64-bit na atomic compare-and-exchange na mga operasyon. Siyempre, kapag lumipat sa 64 bits, bahagyang bababa ang pagganap.
Status ng hash table
Ang bawat pares ng key-value sa isang hash table ay maaaring magkaroon ng isa sa apat na estado:
- Walang laman ang susi at halaga. Sa ganitong estado, ang hash table ay sinisimulan.
- Ang susi ay naisulat na, ngunit ang halaga ay hindi pa naisulat. Kung ang isa pang thread ay kasalukuyang nagbabasa ng data, pagkatapos ay ibabalik itong walang laman. Ito ay normal, ang parehong bagay ay nangyari kung ang isa pang thread ng pagpapatupad ay nagtrabaho nang kaunti nang mas maaga, at pinag-uusapan natin ang tungkol sa isang kasabay na istraktura ng data.
- Parehong ang susi at ang halaga ay naitala.
- Ang halaga ay magagamit sa iba pang mga thread ng pagpapatupad, ngunit ang susi ay hindi pa. Ito ay maaaring mangyari dahil ang modelo ng CUDA programming ay may maluwag na nakaayos na modelo ng memorya. Ito ay normal; sa anumang kaganapan, ang susi ay walang laman pa rin, kahit na ang halaga ay hindi na ganoon.
Ang isang mahalagang nuance ay na kapag ang susi ay naisulat na sa puwang, hindi na ito gumagalaw - kahit na ang susi ay tinanggal, pag-uusapan natin ito sa ibaba.
Gumagana pa nga ang hash table code sa mga maluwag na inayos na mga modelo ng memory kung saan hindi alam ang pagkakasunud-sunod kung saan binasa at isinulat ang memorya. Habang tinitingnan natin ang pagpapasok, paghahanap, at pagtanggal sa isang hash table, tandaan na ang bawat pares ng key-value ay nasa isa sa apat na estado na inilarawan sa itaas.
Pagpasok sa isang hash table
Ang function ng CUDA na naglalagay ng mga pares ng key-value sa isang hash table ay ganito ang hitsura:
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);
}
}
Upang magpasok ng key, umuulit ang code sa hanay ng hash table na nagsisimula sa hash ng inilagay na key. Ang bawat slot sa array ay nagsasagawa ng atomic compare-and-swap operation na nagkukumpara sa key sa slot na iyon sa walang laman. Kung may nakitang mismatch, ang susi sa slot ay ina-update gamit ang nakapasok na key, at pagkatapos ay ibabalik ang orihinal na slot key. Kung ang orihinal na key na ito ay walang laman o tumugma sa ipinasok na key, ang code ay nakakita ng angkop na puwang para sa pagpasok at ipinasok ang inilagay na halaga sa puwang.
Kung sa isang kernel call gpu_hashtable_insert() mayroong maraming mga elemento na may parehong susi, kung gayon ang alinman sa kanilang mga halaga ay maaaring isulat sa key slot. Ito ay itinuturing na normal: ang isa sa mga key-value na pagsusulat sa panahon ng tawag ay magtatagumpay, ngunit dahil ang lahat ng ito ay nangyayari nang magkatulad sa loob ng ilang mga thread ng pagpapatupad, hindi namin mahuhulaan kung aling memory write ang magiging huling isa.
Hash table lookup
Code para sa mga susi sa paghahanap:
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);
}
}
Upang mahanap ang halaga ng isang susi na nakaimbak sa isang talahanayan, umuulit kami sa array na nagsisimula sa hash ng key na hinahanap namin. Sa bawat slot, tinitingnan namin kung ang susi ay ang hinahanap namin, at kung gayon, ibinabalik namin ang halaga nito. Sinusuri din namin kung ang susi ay walang laman, at kung gayon, ibinabagsak namin ang paghahanap.
Kung hindi namin mahanap ang susi, ang code ay nagbabalik ng walang laman na halaga.
Ang lahat ng mga operasyon sa paghahanap na ito ay maaaring isagawa nang sabay-sabay sa pamamagitan ng mga pagpapasok at pagtanggal. Ang bawat pares sa talahanayan ay magkakaroon ng isa sa apat na estado na inilarawan sa itaas para sa daloy.
Tinatanggal sa isang hash table
Code para sa pagtanggal ng mga susi:
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);
}
}
Ang pagtanggal ng susi ay ginagawa sa hindi pangkaraniwang paraan: iniiwan namin ang susi sa talahanayan at minarkahan ang halaga nito (hindi ang susi mismo) bilang walang laman. Ang code na ito ay halos kapareho sa lookup(), maliban na kapag may nakitang tugma sa isang susi, ginagawa nitong walang laman ang halaga nito.
Tulad ng nabanggit sa itaas, kapag ang isang susi ay naisulat sa isang puwang, hindi na ito ililipat. Kahit na ang isang elemento ay tinanggal mula sa talahanayan, ang susi ay nananatili sa lugar, ang halaga nito ay nagiging walang laman. Nangangahulugan ito na hindi namin kailangang gumamit ng atomic write operation para sa halaga ng slot, dahil hindi mahalaga kung ang kasalukuyang halaga ay walang laman o wala - ito ay magiging walang laman.
Pag-resize ng hash table
Maaari mong baguhin ang laki ng hash table sa pamamagitan ng paggawa ng mas malaking table at pagpasok dito ng mga hindi walang laman na elemento mula sa lumang table. Hindi ko ipinatupad ang functionality na ito dahil gusto kong panatilihing simple ang sample code. Bukod dito, sa mga programa ng CUDA, ang paglalaan ng memorya ay madalas na ginagawa sa host code kaysa sa kernel ng CUDA.
Sa artikulo ay naglalarawan kung paano baguhin ang naturang istruktura ng data na protektado ng lock.
pagiging mapagkumpitensya
Sa mga snippet ng code ng function sa itaas gpu_hashtable_insert(), _lookup() ΠΈ _delete() iproseso ang isang pares ng key-value sa isang pagkakataon. At mas mababa gpu_hashtable_insert(), _lookup() ΠΈ _delete() iproseso ang isang hanay ng mga pares nang magkatulad, ang bawat pares sa isang hiwalay na thread ng pagpapatupad ng 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);
}
}
Sinusuportahan ng talahanayan ng hash na lumalaban sa lock ang mga sabay-sabay na pagsingit, paghahanap, at pagtanggal. Dahil ang mga pares ng key-value ay palaging nasa isa sa apat na estado at ang mga susi ay hindi gumagalaw, ginagarantiyahan ng talahanayan ang kawastuhan kahit na ang iba't ibang uri ng mga operasyon ay ginagamit nang sabay-sabay.
Gayunpaman, kung magpoproseso kami ng isang batch ng mga pagpapasok at pagtanggal nang magkatulad, at kung ang input array ng mga pares ay naglalaman ng mga duplicate na key, hindi namin mahuhulaan kung aling mga pares ang "manalo"βay isusulat sa hash table na huling. Sabihin nating tinawag namin ang insertion code na may input array ng mga pares A/0 B/1 A/2 C/3 A/4. Kapag nakumpleto na ang code, magpares B/1 ΠΈ C/3 ay garantisadong naroroon sa talahanayan, ngunit sa parehong oras ang alinman sa mga pares ay lilitaw dito A/0, A/2 o A/4. Ito ay maaaring isang problema o hindi - ang lahat ay nakasalalay sa aplikasyon. Maaaring alam mo nang maaga na walang mga duplicate na key sa input array, o maaaring wala kang pakialam kung aling value ang huling isinulat.
Kung ito ay isang problema para sa iyo, kailangan mong paghiwalayin ang mga duplicate na pares sa iba't ibang CUDA system call. Sa CUDA, ang anumang operasyon na tumatawag sa kernel ay palaging nakumpleto bago ang susunod na kernel call (kahit sa loob ng isang thread. Sa iba't ibang mga thread, ang mga kernel ay pinapagana nang magkatulad). Sa halimbawa sa itaas, kung tumawag ka ng isang kernel na may A/0 B/1 A/2 C/3, at ang isa pa ay may A/4, pagkatapos ay ang susi A makakakuha ng halaga 4.
Ngayon pag-usapan natin kung dapat ang mga function lookup() ΠΈ delete() gumamit ng payak o pabagu-bagong pointer sa isang hanay ng mga pares sa hash table. Isinasaad na:
Maaaring piliin ng compiler na i-optimize ang mga read at writes sa global o shared memory... Maaaring hindi paganahin ang mga optimization na ito gamit ang keyword
volatile: ... anumang reference sa variable na ito ay pinagsama-sama sa isang tunay na memory read o write instruction.
Ang mga pagsasaalang-alang sa kawastuhan ay hindi nangangailangan ng aplikasyon volatile. Kung ang execution thread ay gumagamit ng isang naka-cache na halaga mula sa isang naunang read operation, ito ay gagamit ng bahagyang hindi napapanahong impormasyon. Ngunit gayon pa man, ito ay impormasyon mula sa tamang estado ng hash table sa isang tiyak na sandali ng kernel call. Kung kailangan mong gamitin ang pinakabagong impormasyon, maaari mong gamitin ang index volatile, ngunit pagkatapos ay bahagyang bababa ang pagganap: ayon sa aking mga pagsubok, kapag nagtanggal ng 32 milyong elemento, bumaba ang bilis mula 500 milyong pagtanggal/seg hanggang 450 milyong pagtanggal/seg.
Pagiging Produktibo
Sa pagsubok para sa pagpasok ng 64 milyong elemento at pagtanggal ng 32 milyon sa kanila, kumpetisyon sa pagitan std::unordered_map at halos walang hash table para sa GPU:
std::unordered_map gumugol ng 70 ms sa pagpasok at pag-alis ng mga elemento at pagkatapos ay palayain ang mga ito unordered_map (Ang pag-alis ng milyun-milyong elemento ay nangangailangan ng maraming oras, dahil sa loob unordered_map maramihang mga paglalaan ng memorya ang ginawa). Sa totoo lang, std:unordered_map ganap na magkakaibang mga paghihigpit. Ito ay isang solong CPU thread ng execution, sumusuporta sa mga key-values ββng anumang laki, mahusay na gumaganap sa mataas na rate ng paggamit, at nagpapakita ng matatag na pagganap pagkatapos ng maraming mga pagtanggal.
Ang tagal ng hash table para sa GPU at inter-program na komunikasyon ay 984 ms. Kabilang dito ang oras na ginugol sa paglalagay ng talahanayan sa memorya at pagtanggal nito (paglalaan ng 1 GB ng memory sa isang beses, na tumatagal ng ilang oras sa CUDA), pagpasok at pagtanggal ng mga elemento, at pag-ulit sa mga ito. Ang lahat ng mga kopya papunta at mula sa memorya ng video card ay isinasaalang-alang din.
Ang hash table mismo ay tumagal ng 271 ms upang makumpleto. Kabilang dito ang oras na ginugol ng video card sa pagpasok at pagtanggal ng mga elemento, at hindi isinasaalang-alang ang oras na ginugol sa pagkopya sa memorya at pag-ulit sa resultang talahanayan. Kung ang talahanayan ng GPU ay nabubuhay nang mahabang panahon, o kung ang talahanayan ng hash ay ganap na nakapaloob sa memorya ng video card (halimbawa, upang lumikha ng isang talahanayan ng hash na gagamitin ng ibang GPU code at hindi ang gitnang processor), kung gayon may kaugnayan ang resulta ng pagsusulit.
Ang hash table para sa isang video card ay nagpapakita ng mataas na performance dahil sa mataas na throughput at aktibong parallelization.
Mga hangganan
Ang arkitektura ng hash table ay may ilang isyu na dapat malaman:
- Ang linear probing ay nahahadlangan ng clustering, na nagiging sanhi ng paglalagay ng mga susi sa talahanayan nang hindi gaanong perpekto.
- Ang mga susi ay hindi inaalis gamit ang function
deleteat sa paglipas ng panahon ay kalat nila ang mesa.
Bilang resulta, ang pagganap ng isang hash table ay maaaring unti-unting bumaba, lalo na kung ito ay umiiral nang mahabang panahon at may maraming pagsingit at pagtanggal. Ang isang paraan para mabawasan ang mga kawalan na ito ay ang muling paglalagay sa isang bagong talahanayan na may medyo mababang rate ng paggamit at i-filter ang mga inalis na key sa panahon ng rehashing.
Upang ilarawan ang mga isyung inilarawan, gagamitin ko ang code sa itaas upang lumikha ng isang talahanayan na may 128 milyong elemento at umikot sa 4 milyong elemento hanggang sa mapunan ko ang 124 milyong mga puwang (rate ng paggamit na humigit-kumulang 0,96). Narito ang talahanayan ng resulta, ang bawat hilera ay isang CUDA kernel call upang magpasok ng 4 na milyong bagong elemento sa isang hash table:
Bilis ng paggamit
Tagal ng pagpasok 4 na elemento
0,00
11,608448 ms (361,314798 milyong key/seg.)
0,03
11,751424 ms (356,918799 milyong key/seg.)
0,06
11,942592 ms (351,205515 milyong key/seg.)
0,09
12,081120 ms (347,178429 milyong key/seg.)
0,12
12,242560 ms (342,600233 milyong key/seg.)
0,16
12,396448 ms (338,347235 milyong key/seg.)
0,19
12,533024 ms (334,660176 milyong key/seg.)
0,22
12,703328 ms (330,173626 milyong key/seg.)
0,25
12,884512 ms (325,530693 milyong key/seg.)
0,28
13,033472 ms (321,810182 milyong key/seg.)
0,31
13,239296 ms (316,807174 milyong key/seg.)
0,34
13,392448 ms (313,184256 milyong key/seg.)
0,37
13,624000 ms (307,861434 milyong key/seg.)
0,41
13,875520 ms (302,280855 milyong key/seg.)
0,44
14,126528 ms (296,909756 milyong key/seg.)
0,47
14,399328 ms (291,284699 milyong key/seg.)
0,50
14,690304 ms (285,515123 milyong key/seg.)
0,53
15,039136 ms (278,892623 milyong key/seg.)
0,56
15,478656 ms (270,973402 milyong key/seg.)
0,59
15,985664 ms (262,379092 milyong key/seg.)
0,62
16,668673 ms (251,627968 milyong key/seg.)
0,66
17,587200 ms (238,486174 milyong key/seg.)
0,69
18,690048 ms (224,413765 milyong key/seg.)
0,72
20,278816 ms (206,831789 milyong key/seg.)
0,75
22,545408 ms (186,038058 milyong key/seg.)
0,78
26,053312 ms (160,989275 milyong key/seg.)
0,81
31,895008 ms (131,503463 milyong key/seg.)
0,84
42,103294 ms (99,619378 milyong key/seg.)
0,87
61,849056 ms (67,815164 milyong key/seg.)
0,90
105,695999 ms (39,682713 milyong key/seg.)
0,94
240,204636 ms (17,461378 milyong key/seg.)
Habang tumataas ang paggamit, bumababa ang pagganap. Ito ay hindi kanais-nais sa karamihan ng mga kaso. Kung ang isang application ay nagpasok ng mga elemento sa isang talahanayan at pagkatapos ay itatapon ang mga ito (halimbawa, kapag nagbibilang ng mga salita sa isang libro), hindi ito isang problema. Ngunit kung ang application ay gumagamit ng long-lived hash table (halimbawa, sa isang graphics editor upang mag-imbak ng mga di-bakanteng bahagi ng mga larawan kung saan ang user ay madalas na naglalagay at nagtatanggal ng impormasyon), kung gayon ang pag-uugaling ito ay maaaring maging problema.
At sinukat ang hash table probing depth pagkatapos ng 64 million inserts (utilization factor 0,5). Ang average na depth ay 0,4774, kaya karamihan sa mga key ay nasa pinakamagandang slot o isang slot ang layo mula sa pinakamagandang posisyon. Ang maximum na lalim ng tunog ay 60.
Pagkatapos ay sinukat ko ang lalim ng probing sa isang talahanayan na may 124 milyong pagsingit (utilization factor 0,97). Ang average na lalim ay 10,1757 na, at ang maximum - 6474 (!!). Ang pagganap ng linear sensing ay makabuluhang bumaba sa mataas na mga rate ng paggamit.
Pinakamainam na panatilihing mababa ang rate ng paggamit ng hash table na ito. Ngunit pagkatapos ay pinapataas namin ang pagganap sa gastos ng pagkonsumo ng memorya. Sa kabutihang palad, sa kaso ng 32-bit na mga susi at mga halaga, ito ay maaaring makatwiran. Kung sa halimbawa sa itaas, sa isang talahanayan na may 128 milyong elemento, pinapanatili natin ang kadahilanan ng paggamit na 0,25, kung gayon maaari tayong maglagay ng hindi hihigit sa 32 milyong elemento dito, at ang natitirang 96 milyong mga puwang ay mawawala - 8 byte para sa bawat pares , 768 MB ng nawalang memorya.
Mangyaring tandaan na pinag-uusapan natin ang pagkawala ng memorya ng video card, na isang mas mahalagang mapagkukunan kaysa sa memorya ng system. Bagama't karamihan sa mga modernong desktop graphics card na sumusuporta sa CUDA ay may hindi bababa sa 4 GB ng memorya (sa oras ng pagsulat, ang NVIDIA 2080 Ti ay may 11 GB), hindi pa rin ito ang pinakamatalinong desisyon na mawala ang mga naturang halaga.
Mamaya ay magsusulat ako ng higit pa tungkol sa paglikha ng mga hash table para sa mga video card na walang problema sa lalim ng probing, pati na rin ang mga paraan upang muling gamitin ang mga tinanggal na slot.
Pagsukat ng lalim ng tunog
Upang matukoy ang lalim ng pagsisiyasat ng isang susi, maaari nating kunin ang hash ng susi (ang ideal na index ng talahanayan nito) mula sa aktwal na index ng talahanayan nito:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Dahil sa mahika ng mga binary na numero ng complement ng dalawa at ang katotohanan na ang kapasidad ng hash table ay dalawa sa kapangyarihan ng dalawa, gagana ang diskarteng ito kahit na inilipat ang key index sa simula ng talahanayan. Kunin natin ang isang susi na na-hash sa 1, ngunit ipinasok sa slot 3. Pagkatapos, para sa isang mesa na may kapasidad na 4 ay nakukuha natin (3 β 1) & 3, na katumbas ng 2.
Konklusyon
Kung mayroon kang mga katanungan o komento, mangyaring mag-email sa akin sa o magbukas ng bagong paksa sa.
Ang code na ito ay isinulat sa ilalim ng inspirasyon mula sa mahusay na mga artikulo:
Sa hinaharap, patuloy akong magsusulat tungkol sa mga pagpapatupad ng hash table para sa mga video card at susuriin ang kanilang pagganap. Kasama sa mga plano ko ang chaining, Robin Hood hashing, at cuckoo hashing gamit ang atomic operations sa mga istruktura ng data na GPU friendly.
Pinagmulan: www.habr.com