Kuring dipasang dina Github .
Ieu tabel hash GPU basajan sanggup ngolah ratusan juta inserts per detik. Dina laptop NVIDIA GTX 1060 kuring, kodena nyelapkeun 64 juta pasangan nilai konci anu dihasilkeun sacara acak sakitar 210 ms sareng ngahapus 32 juta pasang dina sakitar 64 ms.
Hartina, laju dina laptop kurang leuwih 300 juta inserts/detik jeung 500 juta deletes/detik.
tabél ieu ditulis dina CUDA, sanajan téhnik sarua bisa dilarapkeun ka HLSL atanapi GLSL. Palaksanaanna ngagaduhan sababaraha watesan pikeun mastikeun kinerja anu luhur dina kartu vidéo:
- Ngan konci 32-bit sareng nilai anu sami anu diolah.
- Méja hash gaduh ukuran tetep.
- Sareng ukuran ieu kedah sami sareng dua kakuatan.
Pikeun konci na nilai, Anjeun kudu cagar spidol delimiter basajan (dina kode luhur ieu 0xffffffff).
Méja Hash tanpa konci
Tabel hash ngagunakeun alamat kabuka kalayan , nyaeta, ngan saukur hiji Asép Sunandar Sunarya ti pasangan konci-nilai nu disimpen dina mémori jeung boga kinerja cache unggul. Sami teu tiasa nyarios chaining, anu ngalibatkeun milarian pointer dina daptar numbu. Méja hash mangrupikeun saderhana pikeun nyimpen unsur KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Ukuran tabel nyaéta kakuatan dua, lain nomer prima, sabab hiji instruksi gancang cukup pikeun nerapkeun pow2 / AND topeng, tapi operator modulus loba laun. Ieu penting dina kasus probing linier, saprak dina tabel linier lookup indéks slot kudu dibungkus dina unggal slot. Jeung salaku hasilna, biaya operasi ditambahkeun modulo dina unggal slot.
tabél ngan nyimpen konci na nilai pikeun tiap unsur, teu hash sahiji konci. Kusabab tabél ngan nyimpen konci 32-bit, Hash diitung gancang pisan. Kodeu di luhur ngagunakeun Hash Murmur3, nu ukur ngalakukeun sababaraha shifts, XORs na multiplications.
Méja Hash ngagunakeun téknik panyalindungan ngonci anu bebas tina urutan memori. Sanaos sababaraha operasi nyerat ngaganggu tatanan operasi sapertos kitu, tabel hash tetep bakal ngajaga kaayaan anu leres. Urang bakal ngobrol ngeunaan ieu di handap. Téhnikna tiasa dianggo saé sareng kartu vidéo anu ngajalankeun rébuan benang sakaligus.
Konci sareng nilai dina tabel hash diinisialisasi janten kosong.
Kodeu tiasa dirobih pikeun nanganan konci sareng nilai 64-bit ogé. Konci merlukeun operasi atom maca, nulis, jeung ngabandingkeun-jeung-swap. Sareng nilai ngabutuhkeun operasi maca sareng nyerat atom. Untungna, dina CUDA, operasi baca-tulis pikeun nilai 32- sareng 64-bit nyaéta atom salami aranjeunna dijajarkeun sacara alami (tingali di handap). ), sareng kartu vidéo modern ngadukung operasi ngabandingkeun-sareng-tukar atom 64-bit. Tangtosna, nalika pindah ka 64 bit, kinerja bakal rada turun.
kaayaan tabel Hash
Unggal pasangan konci-nilai dina tabel hash tiasa gaduh salah sahiji tina opat nagara bagian:
- Konci sareng nilai kosong. Dina kaayaan ieu, tabel hash diinisialisasi.
- Koncina parantos dituliskeun, tapi nilaina henteu acan ditulis. Lamun thread sejen ayeuna maca data, lajeng mulih kosong. Ieu normal, hal anu sarua bakal kajadian lamun thread sejen tina palaksanaan geus digawé saeutik saméméhna, sarta kami ngawangkong ngeunaan struktur data babarengan.
- Boh konci sareng nilaina dirékam.
- Nilaina sayogi pikeun utas palaksanaan anu sanés, tapi koncina henteu acan. Ieu tiasa kajantenan kusabab modél program CUDA gaduh modél mémori anu diurutkeun sacara bébas. Ieu normal; dina kaayaan naon waé, koncina masih kosong, sanaos nilaina henteu deui.
Hiji nuansa penting éta sakali konci geus ditulis kana slot, éta euweuh gerak - sanajan konci dihapus, urang bakal ngobrol ngeunaan ieu di handap.
Kodeu tabel hash malah tiasa dianggo sareng modél mémori anu dipesen sacara bébas dimana urutan anu dibaca sareng ditulis mémori henteu dipikanyaho. Nalika urang ningali sisipan, milarian, sareng ngahapus dina tabel hash, émut yén unggal pasangan konci-nilai aya dina salah sahiji tina opat nagara bagian anu dijelaskeun di luhur.
Nyelapkeun kana méja hash
Fungsi CUDA anu nyelapkeun pasangan konci-nilai kana tabel hash sapertos kieu:
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);
}
}
Pikeun nyelapkeun konci hiji, kode iterates ngaliwatan Asép Sunandar Sunarya tabel hash dimimitian ku hash sahiji konci diselapkeun. Unggal slot dina Asép Sunandar Sunarya ngalakukeun operasi ngabandingkeun-na-swap atom nu compares konci dina slot éta kosong. Mun mismatch a dideteksi, konci dina slot diropéa kalawan konci diselapkeun, lajeng konci slot aslina balik. Upami konci aslina ieu kosong atanapi cocog sareng konci anu diselapkeun, kode éta mendakan slot anu cocog pikeun diselapkeun sareng diselapkeun nilai anu diselapkeun kana slot.
Lamun dina hiji panggero kernel gpu_hashtable_insert() Aya sababaraha elemen sareng konci anu sami, teras salah sahiji nilaina tiasa diserat kana slot konci. Ieu dianggap normal: salah sahiji konci-nilai nulis salila nelepon bakal sukses, tapi saprak kabeh ieu lumangsung dina paralel dina sababaraha threads palaksanaan, urang teu bisa ngaduga mana nulis memori bakal jadi panungtungan.
Hash méja lookup
Kodeu pikeun milarian konci:
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);
}
}
Pikeun manggihan nilai konci disimpen dina tabél, urang iterate ngaliwatan Asép Sunandar Sunarya dimimitian ku hash sahiji konci kami pilari. Dina unggal slot, urang pariksa naha konci anu urang pilari, sarta lamun kitu, urang balik nilai na. Urang ogé pariksa naha konci éta kosong, sarta lamun kitu, urang abort pilarian.
Lamun urang teu bisa manggihan konci, kode mulih hiji nilai kosong.
Sadaya operasi pamilarian ieu tiasa dilakukeun sakaligus ngalangkungan sisipan sareng ngahapus. Unggal pasangan dina tabél bakal boga salah sahiji opat nagara bagian ditétélakeun di luhur pikeun aliran.
Ngahapus dina tabel hash
Kode pikeun mupus konci:
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);
}
}
Ngahapus konci dilakukeun ku cara anu teu biasa: urang ngantepkeun konci dina tabél sareng nandaan nilaina (sanés koncina sorangan) kosong. Kode ieu pisan sarupa lookup(), iwal mun cocok kapanggih dina konci, ngajadikeun nilai na kosong.
Sakumaha didadarkeun di luhur, sakali hiji konci ditulis kana slot, eta geus euweuh dipindahkeun. Malah lamun hiji unsur dihapus tina tabél, konci tetep dina tempatna, nilaina saukur jadi kosong. Ieu ngandung harti yén urang teu kedah nganggo operasi nulis atom pikeun nilai slot , sabab henteu masalah naha nilai ayeuna kosong atanapi henteu - eta tetep bakal kosong.
Ngarobah ukuran hiji méja hash
Anjeun tiasa ngarobah ukuran tabel Hash ku nyieun hiji méja gedé tur nyelapkeun elemen non-kosong tina tabel heubeul ka dinya. Kuring henteu nerapkeun fungsionalitas ieu kusabab kuring hoyong tetep kode sampelna basajan. Leuwih ti éta, dina program CUDA, alokasi memori mindeng dipigawé dina kode host tinimbang dina kernel CUDA.
Tulisanana ngajelaskeun kumaha carana ngarobih struktur data anu ditangtayungan ku konci sapertos kitu.
Daya saing
Dina snippét kode fungsi luhur gpu_hashtable_insert(), _lookup() и _delete() ngolah hiji pasangan konci-nilai dina hiji waktu. Jeung handap gpu_hashtable_insert(), _lookup() и _delete() ngolah susunan pasangan paralel, unggal pasangan dina thread palaksanaan GPU misah:
// 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);
}
}
Méja hash anu tahan konci ngadukung sisipan, panéangan, sareng ngahapus sakaligus. Kusabab pasangan konci-nilai sok aya dina salah sahiji opat nagara bagian jeung kenop teu mindahkeun, tabél ngajamin correctness sanajan tipena béda operasi dipaké sakaligus.
Sanajan kitu, lamun urang ngolah bets insertions na ngahapus dina paralel, sarta lamun Asép Sunandar Sunarya input pasangan ngandung duplikat konci, lajeng urang moal bisa ngaduga pasangan nu bakal "meunang" - bakal ditulis kana tabel Hash panungtungan. Hayu urang sebutkeun kode sisipan kalawan Asép Sunandar Sunarya input pasangan A/0 B/1 A/2 C/3 A/4. Nalika kode réngsé, pasang B/1 и C/3 dijamin bakal hadir dina tabél, tapi dina waktos anu sareng salah sahiji pasangan bakal muncul di dinya A/0, A/2 atawa A/4. Ieu tiasa atanapi henteu janten masalah - sadayana gumantung kana aplikasina. Anjeun bisa jadi terang sateuacanna yén euweuh duplikat konci dina Asép Sunandar Sunarya input, atawa anjeun bisa jadi teu paduli nu nilai ieu ditulis panungtungan.
Upami ieu mangrupikeun masalah pikeun anjeun, maka anjeun kedah misahkeun pasangan duplikat kana telepon sistem CUDA anu béda. Dina CUDA, sagala operasi nu nelepon kernel salawasna ngalengkepan saméméh panggero kernel salajengna (sahenteuna dina hiji thread. Dina threads béda, kernels dieksekusi paralel). Dina conto di luhur, lamun nelepon hiji kernel kalawan A/0 B/1 A/2 C/3, jeung lianna kalayan A/4, lajeng konci A bakal meunang nilai 4.
Ayeuna hayu urang ngobrol ngeunaan naha fungsi kedah lookup() и delete() ngagunakeun pointer polos atawa volatile ka Asép Sunandar Sunarya ti pasangan dina tabel Hash. nyatakeun yén:
Kompiler tiasa milih pikeun ngaoptimalkeun maca sareng nyerat kana mémori global atanapi dibagikeun... Optimasi ieu tiasa ditumpurkeun nganggo kecap konci
volatile: ... sagala rujukan pikeun variabel ieu disusun kana memori nyata maca atawa nulis instruksi.
Pertimbangan correctness teu merlukeun aplikasi volatile. Upami utas palaksanaan nganggo nilai sindangan tina operasi anu dibaca sateuacana, maka éta bakal ngagunakeun inpormasi anu rada luntur. Tapi tetep, ieu mangrupikeun inpormasi tina kaayaan anu leres dina tabel hash dina waktos anu tangtu dina sauran kernel. Lamun perlu ngagunakeun inpo panganyarna, Anjeun tiasa make indéks dina volatile, tapi lajeng kinerja bakal rada ngurangan: nurutkeun tés kuring, nalika ngahapus 32 juta elemen, laju turun tina 500 juta ngahapus / detik ka 450 juta ngahapus / detik.
kakuwatan keur ngasilkeun
Dina uji pikeun inserting 64 juta elemen sarta mupus 32 juta di antarana, kompetisi antara std::unordered_map sareng ampir teu aya tabel hash pikeun GPU:
std::unordered_map spent 70 mdet inserting jeung nyoplokkeun elemen lajeng freeing aranjeunna unordered_map (Nyingkirkeun jutaan elemen butuh seueur waktos, sabab di jero unordered_map sababaraha alokasi memori dijieun). Jujur ngomong, std:unordered_map larangan lengkep beda. Ieu mangrupikeun utas palaksanaan CPU tunggal, ngadukung nilai konci tina ukuran naon waé, ngalaksanakeun saé dina tingkat panggunaan anu luhur, sareng nunjukkeun prestasi anu stabil saatos sababaraha ngahapus.
Durasi tabel hash pikeun GPU sareng komunikasi antar-program éta 984 mdet. Ieu ngawengku waktu spent nempatkeun tabel dina mémori jeung mupus eta (allocating 1 GB memori hiji waktu, nu butuh sababaraha waktos di CUDA), inserting sarta mupus elemen, sarta iterating leuwih aranjeunna. Sadaya salinan ka sareng ti mémori kartu vidéo ogé dipertimbangkeun.
Méja hash sorangan nyandak 271 mdet pikeun réngsé. Ieu ngawengku waktu spent ku kartu vidéo inserting sarta mupus elemen, sarta teu tumut kana akun waktos spent nyalin kana memori tur iterating leuwih tabel hasilna. Lamun tabel GPU hirup lila, atawa lamun tabel hash dikandung sapinuhna dina mémori kartu vidéo (contona, nyieun tabel hash nu bakal dipaké ku kode GPU sejen tur teu prosesor sentral), lajeng. hasil tés relevan.
Méja hash pikeun kartu vidéo nunjukkeun kinerja anu luhur kusabab throughput anu luhur sareng paralélisasi aktip.
shortcomings
Arsitéktur tabel hash ngagaduhan sababaraha masalah anu kedah diperhatoskeun:
- probing linier ieu hampered ku clustering, nu ngabalukarkeun kenop dina tabél ditempatkeun kirang ti sampurna.
- Konci henteu dipiceun nganggo fungsina
deletesarta leuwih waktos aranjeunna clutter méja.
Hasilna, kinerja tabel hash laun-laun tiasa nguraikeun, khususna upami parantos lami sareng seueur sisipan sareng ngahapus. Salah sahiji cara pikeun ngirangan kalemahan ieu nyaéta ngarobih deui kana méja énggal kalayan tingkat pamakean anu cukup rendah sareng nyaring konci anu dipiceun salami rehashing.
Pikeun ngagambarkeun masalah anu dijelaskeun, kuring bakal nganggo kode di luhur pikeun nyiptakeun méja kalayan 128 juta elemen sareng puteran 4 juta elemen dugi ka kuring ngeusian 124 juta liang (laju utilization ngeunaan 0,96). Ieu tabel hasil, unggal baris nyaéta panggero kernel CUDA pikeun nyelapkeun 4 juta elemen anyar kana hiji méja hash:
Laju pamakéan
Durasi sisipan 4 elemen
0,00
11,608448 mdet (361,314798 juta konci/detik)
0,03
11,751424 mdet (356,918799 juta konci/detik)
0,06
11,942592 mdet (351,205515 juta konci/detik)
0,09
12,081120 mdet (347,178429 juta konci/detik)
0,12
12,242560 mdet (342,600233 juta konci/detik)
0,16
12,396448 mdet (338,347235 juta konci/detik)
0,19
12,533024 mdet (334,660176 juta konci/detik)
0,22
12,703328 mdet (330,173626 juta konci/detik)
0,25
12,884512 mdet (325,530693 juta konci/detik)
0,28
13,033472 mdet (321,810182 juta konci/detik)
0,31
13,239296 mdet (316,807174 juta konci/detik)
0,34
13,392448 mdet (313,184256 juta konci/detik)
0,37
13,624000 mdet (307,861434 juta konci/detik)
0,41
13,875520 mdet (302,280855 juta konci/detik)
0,44
14,126528 mdet (296,909756 juta konci/detik)
0,47
14,399328 mdet (291,284699 juta konci/detik)
0,50
14,690304 mdet (285,515123 juta konci/detik)
0,53
15,039136 mdet (278,892623 juta konci/detik)
0,56
15,478656 mdet (270,973402 juta konci/detik)
0,59
15,985664 mdet (262,379092 juta konci/detik)
0,62
16,668673 mdet (251,627968 juta konci/detik)
0,66
17,587200 mdet (238,486174 juta konci/detik)
0,69
18,690048 mdet (224,413765 juta konci/detik)
0,72
20,278816 mdet (206,831789 juta konci/detik)
0,75
22,545408 mdet (186,038058 juta konci/detik)
0,78
26,053312 mdet (160,989275 juta konci/detik)
0,81
31,895008 mdet (131,503463 juta konci/detik)
0,84
42,103294 mdet (99,619378 juta konci/detik)
0,87
61,849056 mdet (67,815164 juta konci/detik)
0,90
105,695999 mdet (39,682713 juta konci/detik)
0,94
240,204636 mdet (17,461378 juta konci/detik)
Nalika utilization ningkat, kinerja turun. Ieu teu desirable dina kalolobaan kasus. Upami aplikasi nyelapkeun elemen kana méja teras miceunana (contona, nalika ngitung kecap dina buku), maka ieu sanés masalah. Tapi lamun aplikasi ngagunakeun tabel hash lila-cicing (contona, dina redaktur grafik pikeun nyimpen bagian non-kosong gambar dimana pamaké remen inserts sarta mupus informasi), lajeng kabiasaan ieu bisa jadi masalah.
Sarta diukur dina tabel hash probing jero sanggeus 64 juta inserts (faktor utilization 0,5). Jero rata éta 0,4774, jadi paling konci éta boh dina pangalusna mungkin slot atawa hiji slot jauh ti posisi pangalusna. Jero sora maksimum nyaéta 60.
Kuring lajeng ngukur jero probing dina tabel kalawan 124 juta inserts (faktor utilization 0,97). Jero rata-rata parantos 10,1757, sareng maksimum - 6474 (!!). Kinerja sensing linier turun sacara signifikan dina tingkat panggunaan anu luhur.
Hadé pisan mun éta tetep laju utilization tabel hash ieu low. Tapi lajeng urang ningkatkeun kinerja di expense konsumsi memori. Untungna, dina kasus konci sareng nilai 32-bit, ieu tiasa diyakinkeun. Lamun dina conto di luhur, dina tabel kalawan 128 juta elemen, urang tetep faktor utilization 0,25, teras urang tiasa nempatkeun euweuh leuwih ti 32 juta elemen di dinya, sarta sésana 96 juta liang bakal leungit - 8 bait pikeun tiap pasangan. , 768 MB memori leungit.
Punten dicatet yén urang ngobrol ngeunaan leungitna mémori kartu vidéo, anu mangrupikeun sumber anu langkung berharga tibatan mémori sistem. Sanaos kalolobaan kartu grafik desktop modéren anu ngadukung CUDA gaduh mémori sahenteuna 4 GB (dina waktos nyerat, NVIDIA 2080 Ti gaduh 11 GB), éta tetep henteu janten kaputusan anu paling bijaksana pikeun kaleungitan jumlah sapertos kitu.
Engké kuring bakal nulis leuwih seueur tentang nyieun tabel Hash pikeun kartu vidéo nu teu boga masalah sareng jero probing, kitu ogé cara maké deui slot dihapus.
Pangukuran jero sora
Pikeun nangtukeun jero probing konci, urang tiasa nimba hash konci (indéks tabel idéal na) tina indéks tabel sabenerna na:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Kusabab magic tina dua dua urang pelengkap angka binér jeung kanyataan yén kapasitas tabel Hash nyaeta dua kakuatan dua, pendekatan ieu bakal dianggo sanajan indéks konci dipindahkeun ka awal tabél. Hayu urang nyandak konci anu hashed ka 1, tapi diselapkeun kana slot 3. Lajeng pikeun tabel kalawan kapasitas 4 urang meunang. (3 — 1) & 3, anu sarua jeung 2.
kacindekan
Upami Anjeun gaduh patarosan atanapi komentar, mangga surélék kuring di atawa muka topik anyar dina.
Kode ieu ditulis dina inspirasi tina artikel alus teuing:
Dina mangsa nu bakal datang, kuring bakal neruskeun nulis ngeunaan palaksanaan tabel Hash pikeun kartu vidéo jeung nganalisis kinerja maranéhanana. Rencana abdi kalebet chaining, Robin Hood hashing, sareng cuckoo hashing ngagunakeun operasi atom dina struktur data anu ramah GPU.
sumber: www.habr.com