Aku dikirim ing Github .
Iki minangka tabel hash GPU sing prasaja sing bisa ngolah atusan yuta sisipan per detik. Ing laptop NVIDIA GTX 1060, kode kasebut nglebokake 64 yuta pasangan nilai kunci sing digawe kanthi acak ing udakara 210 ms lan mbusak 32 yuta pasangan ing udakara 64 ms.
Tegese, kacepetan ing laptop kira-kira 300 yuta sisipan / detik lan 500 yuta mbusak / detik.
Tabel kasebut ditulis ing CUDA, sanajan teknik sing padha bisa ditrapake kanggo HLSL utawa GLSL. Implementasine duwe sawetara watesan kanggo njamin kinerja dhuwur ing kertu video:
- Mung tombol 32-bit lan nilai sing padha diproses.
- Tabel hash nduweni ukuran tetep.
- Lan ukuran iki kudu padha karo loro kanggo daya.
Kanggo tombol lan nilai, sampeyan kudu leladen panandha delimiter prasaja (ing kode ndhuwur iki 0xffffffff).
Meja hash tanpa kunci
Tabel hash nggunakake alamat mbukak karo , yaiku, iku mung array saka pasangan kunci-nilai sing disimpen ing memori lan nduweni kinerja cache unggul. Padha ora bisa ngandika kanggo chaining, kang melu nggoleki pointer ing dhaftar disambung. Tabel hash minangka unsur nyimpen array sing prasaja KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Ukuran meja punika daya saka loro, ora nomer prima, amarga siji instruksi cepet cukup kanggo aplikasi pow2 / topeng AND, nanging operator modulus luwih alon. Iki penting ing kasus linear probing, wiwit ing tabel linear lookup indeks slot kudu kebungkus ing saben slot . Lan minangka asil, biaya operasi ditambahake modulo ing saben slot .
Tabel mung nyimpen tombol lan nilai kanggo saben unsur, ora hash saka tombol. Wiwit tabel mung nyimpen tombol 32-bit, hash diitung kanthi cepet. Kode ing ndhuwur nggunakake hash Murmur3, sing mung nindakake sawetara shift, XORs lan multiplications.
Tabel hash nggunakake teknik proteksi ngunci sing ora gumantung saka urutan memori. Sanajan sawetara operasi nulis ngganggu urutan operasi liyane, tabel hash isih bakal njaga kahanan sing bener. Kita bakal ngomong babagan iki ing ngisor iki. Teknik kasebut bisa digunakake kanthi apik karo kertu video sing mbukak ewonan utas bebarengan.
Tombol lan nilai ing tabel hash diinisialisasi dadi kosong.
Kode kasebut bisa diowahi kanggo nangani kunci lan nilai 64-bit uga. Tombol mbutuhake operasi maca, nulis, lan mbandhingake-lan-swap atom. Lan nilai mbutuhake operasi maca lan nulis atom. Untunge, ing CUDA, operasi maca-tulis kanggo nilai 32- lan 64-bit iku atom anggere padha selaras kanthi alami (ndeleng ngisor). ), lan kertu video modern ndhukung operasi banding-lan-ijol-ijolan atom 64-bit. Mesthi, nalika pindhah menyang 64 bit, kinerja bakal rada suda.
Status tabel hash
Saben pasangan nilai kunci ing tabel hash bisa duwe salah siji saka patang negara:
- Tombol lan nilai kosong. Ing negara iki, tabel hash diwiwiti.
- Tombol wis ditulis, nanging regane durung ditulis. Yen thread liyane lagi maca data, banjur bali kosong. Iki normal, bab sing padha bakal kedaden yen thread liyane saka eksekusi wis makarya sethitik sadurungé, lan kita ngomong bab struktur data bebarengan.
- Loro-lorone tombol lan nilai direkam.
- Nilai kasedhiya kanggo thread eksekusi liyane, nanging kuncine durung. Iki bisa kedadeyan amarga model pemrograman CUDA duwe model memori sing diurutake. Iki normal; ing acara apa wae, kunci isih kosong, sanajan regane ora ana maneh.
Nuance penting iku yen tombol wis ditulis kanggo slot , iku ora maneh gerakane - malah yen tombol wis dibusak, kita bakal pirembagan bab iki ing ngisor iki.
Kode tabel hash malah bisa digunakake karo model memori sing diurutake kanthi urutan memori sing diwaca lan ditulis ora dingerteni. Nalika ndeleng sisipan, goleki, lan pambusakan ing tabel hash, elinga yen saben pasangan nilai kunci ana ing salah siji saka patang negara sing kasebut ing ndhuwur.
Nglebokake menyang tabel hash
Fungsi CUDA sing nglebokake pasangan kunci-nilai menyang tabel hash katon kaya iki:
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);
}
}
Kanggo nglebokake kunci, kode kasebut diulang liwat array tabel hash sing diwiwiti karo hash tombol sing dilebokake. Saben slot ing array nindakake operasi mbandhingake-lan-swap atom sing mbandhingake kunci ing slot kasebut dadi kosong. Yen mismatch dideteksi, tombol ing slot dianyari karo tombol dipasang, banjur tombol slot asli bali. Yen kunci asli iki kosong utawa cocog karo tombol sing dipasang, kode kasebut nemokake slot sing cocog kanggo dilebokake lan nglebokake nilai sing dipasang ing slot kasebut.
Yen ing siji telpon kernel gpu_hashtable_insert() ana pirang-pirang unsur kanthi tombol sing padha, mula nilai-nilai kasebut bisa ditulis ing slot tombol. Iki dianggep normal: salah siji saka tombol-nilai nulis sak telpon bakal kasil, nanging amarga kabeh iki kedaden ing podo karo ing sawetara Utas saka eksekusi, kita ora bisa prédhiksi kang nulis memori bakal pungkasan.
Goleki tabel hash
Kode kanggo nggoleki kunci:
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);
}
}
Kanggo nemokake nilai tombol sing disimpen ing tabel, kita ngulang array sing diwiwiti kanthi hash tombol sing kita goleki. Ing saben slot, kita mriksa yen tombol iku siji kita looking for, lan yen mangkono, kita bali regane. Kita uga mriksa yen tombol kosong, lan yen mangkono, kita mbatalake panelusuran.
Yen kita ora bisa nemokake kunci, kode kasebut ngasilake nilai kosong.
Kabeh operasi telusuran iki bisa ditindakake bebarengan liwat sisipan lan pambusakan. Saben pasangan ing tabel bakal duwe salah siji saka papat negara sing diterangake ing ndhuwur kanggo aliran.
Mbusak ing tabel hash
Kode kanggo mbusak kunci:
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);
}
}
Mbusak tombol ditindakake kanthi cara sing ora biasa: kita ninggalake tombol ing meja lan menehi tandha nilai (dudu tombol kasebut) kosong. Kode iki meh padha karo lookup(), kajaba yen cocog ditemokaké ing tombol, iku ndadekake regane kosong.
Kaya kasebut ing ndhuwur, yen tombol wis ditulis kanggo slot , iku ora maneh dipindhah. Sanajan ana unsur dibusak saka meja, tombol tetep ing panggonan, regane mung dadi kosong. Iki tegese kita ora perlu nggunakake operasi nulis atom kanggo Nilai slot , amarga iku ora Matter apa nilai saiki kosong utawa ora - isih bakal kosong.
Ngowahi ukuran tabel hash
Sampeyan bisa ngganti ukuran tabel hash kanthi nggawe tabel sing luwih gedhe lan nglebokake unsur non-kosong saka meja lawas menyang. Aku ora ngleksanakake fungsi iki amarga aku pengin tetep kode sampel prasaja. Kajaba iku, ing program CUDA, alokasi memori asring ditindakake ing kode host tinimbang ing kernel CUDA.
Ing artikel nerangake carane ngowahi struktur data sing dilindhungi kunci kasebut.
Daya saing
Ing cuplikan kode fungsi ing ndhuwur gpu_hashtable_insert(), _lookup() и _delete() proses siji pasangan kunci-nilai sekaligus. Lan ngisor gpu_hashtable_insert(), _lookup() и _delete() proses array pasangan kanthi podo karo, saben pasangan ing utas eksekusi GPU sing kapisah:
// 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);
}
}
Tabel hash tahan kunci ndhukung sisipan, telusuran, lan pambusakan bebarengan. Amarga pasangan kunci-nilai tansah ing salah siji saka patang negara lan tombol ora pindhah, Tabel njamin bener sanajan macem-macem jinis operasi digunakake bebarengan.
Nanging, yen kita proses kumpulan sisipan lan pambusakan ing podo karo, lan yen Uploaded input saka pasangan ngandhut duplikat tombol, banjur kita ora bisa kanggo prédhiksi kang pasangan bakal "menang" -bakal ditulis ing tabel hash pungkasan. Ayo kita sebut kode sisipan kanthi pasangan input A/0 B/1 A/2 C/3 A/4. Nalika kode rampung, pasangan B/1 и C/3 dijamin bakal ana ing meja, nanging ing wektu sing padha, pasangan bakal katon ing A/0, A/2 utawa A/4. Iki bisa dadi masalah utawa ora - kabeh gumantung ing aplikasi. Sampeyan bisa uga ngerti ing advance sing ora ana duplikat tombol ing Uploaded input, utawa sampeyan bisa uga ora Care kang Nilai ditulis pungkasan.
Yen iki masalah kanggo sampeyan, sampeyan kudu misahake pasangan duplikat menyang telpon sistem CUDA beda. Ing CUDA, operasi apa wae sing nelpon kernel mesthi rampung sadurunge telpon kernel sabanjure (paling ora ing siji utas. Ing benang sing beda-beda, kernel dieksekusi kanthi paralel). Ing conto ing ndhuwur, yen sampeyan nelpon siji kernel karo A/0 B/1 A/2 C/3, lan liyane karo A/4, banjur kunci A bakal entuk nilai 4.
Saiki ayo ngomong babagan apa fungsi kudu lookup() и delete() nggunakake pointer kosong utawa molah malih menyang Uploaded saka pasangan ing tabel hash. nyatakake yen:
Compiler bisa milih kanggo ngoptimalake maca lan nulis menyang memori global utawa bareng ... Optimizations iki bisa dipateni nggunakake tembung kunci
volatile: ... sembarang referensi kanggo variabel iki nyawiji menyang memori nyata maca utawa nulis instruction.
Pertimbangan sing bener ora mbutuhake aplikasi volatile. Yen thread eksekusi nggunakake nilai cache saka operasi maca sadurungé, banjur bakal nggunakake informasi rada outdated. Nanging isih, iki informasi saka negara bener saka tabel hash ing wayahe tartamtu saka telpon kernel. Yen sampeyan kudu nggunakake informasi paling anyar, sampeyan bisa nggunakake indeks volatile, nanging banjur kinerja bakal rada suda: miturut tesku, nalika mbusak 32 yuta unsur, kacepetan mudhun saka 500 yuta pambusakan / detik dadi 450 yuta pambusakan / detik.
Produktivitas
Ing test kanggo masang 64 yuta unsur lan mbusak 32 yuta mau, kompetisi antarane std::unordered_map lan meh ora ana tabel hash kanggo GPU:
std::unordered_map ngginakaken 70 ms nglebokake lan mbusak unsur banjur mbebasake unordered_map (nyingkirake mayuta-yuta unsur mbutuhake wektu akeh, amarga ing njero unordered_map sawetara alokasi memori digawe). Jujur ngomong, std:unordered_map watesan temen beda. Iki minangka utas eksekusi CPU tunggal, ndhukung nilai kunci saka ukuran apa wae, nindakake kanthi apik ing tingkat panggunaan sing dhuwur, lan nuduhake kinerja sing stabil sawise pirang-pirang pambusakan.
Durasi tabel hash kanggo komunikasi GPU lan antar program yaiku 984 ms. Iki kalebu wektu ngginakaken manggonke meja ing memori lan mbusak (nyedhiyakake 1 GB memori siji wektu, kang njupuk sawetara wektu ing CUDA), nglebokake lan mbusak unsur, lan iterasi liwat mau. Kabeh salinan menyang lan saka memori kertu video uga dianggep.
Tabel hash dhewe njupuk 271 ms kanggo ngrampungake. Iki kalebu wektu ngginakaken dening kertu video nglebokake lan mbusak unsur, lan ora njupuk menyang akun wektu ngginakaken Nyalin menyang memori lan iterating liwat Tabel asil. Yen tabel GPU urip kanggo dangu, utawa yen tabel hash ana ing memori kertu video (contone, kanggo nggawe tabel hash sing bakal digunakake dening kode GPU liyane lan ora prosesor tengah), banjur asil test cocog.
Tabel hash kanggo kertu video nuduhake kinerja dhuwur amarga throughput dhuwur lan paralelisasi aktif.
Kekurangan
Arsitektur tabel hash duwe sawetara masalah sing kudu dingerteni:
- Penyelidikan linear diganggu dening clustering, sing nyebabake tombol ing meja diselehake kurang saka sampurna.
- Tombol ora dibusak nggunakake fungsi
deletelan liwat wektu padha clutter meja.
Akibaté, kinerja tabel hash bisa mboko sithik mudhun, utamané yen ana kanggo dangu lan wis akeh sisipan lan mbusak. Salah sawijining cara kanggo nyuda kekurangan kasebut yaiku rehash menyang tabel anyar kanthi tingkat panggunaan sing cukup sithik lan nyaring tombol sing dicopot sajrone rehashing.
Kanggo ilustrasi masalah diterangake, Aku bakal nggunakake kode ing ndhuwur kanggo nggawe meja karo 128 yuta unsur lan daur ulang liwat 4 yuta unsur nganti aku wis kapenuhan 124 yuta slot (tingkat panggunaan bab 0,96). Iki minangka tabel asil, saben baris minangka panggilan kernel CUDA kanggo nglebokake 4 yuta unsur anyar menyang siji tabel hash:
Tingkat panggunaan
Durasi nglebokake 4 unsur
0,00
11,608448 ms (361,314798 yuta tombol/detik)
0,03
11,751424 ms (356,918799 yuta tombol/detik)
0,06
11,942592 ms (351,205515 yuta tombol/detik)
0,09
12,081120 ms (347,178429 yuta tombol/detik)
0,12
12,242560 ms (342,600233 yuta tombol/detik)
0,16
12,396448 ms (338,347235 yuta tombol/detik)
0,19
12,533024 ms (334,660176 yuta tombol/detik)
0,22
12,703328 ms (330,173626 yuta tombol/detik)
0,25
12,884512 ms (325,530693 yuta tombol/detik)
0,28
13,033472 ms (321,810182 yuta tombol/detik)
0,31
13,239296 ms (316,807174 yuta tombol/detik)
0,34
13,392448 ms (313,184256 yuta tombol/detik)
0,37
13,624000 ms (307,861434 yuta tombol/detik)
0,41
13,875520 ms (302,280855 yuta tombol/detik)
0,44
14,126528 ms (296,909756 yuta tombol/detik)
0,47
14,399328 ms (291,284699 yuta tombol/detik)
0,50
14,690304 ms (285,515123 yuta tombol/detik)
0,53
15,039136 ms (278,892623 yuta tombol/detik)
0,56
15,478656 ms (270,973402 yuta tombol/detik)
0,59
15,985664 ms (262,379092 yuta tombol/detik)
0,62
16,668673 ms (251,627968 yuta tombol/detik)
0,66
17,587200 ms (238,486174 yuta tombol/detik)
0,69
18,690048 ms (224,413765 yuta tombol/detik)
0,72
20,278816 ms (206,831789 yuta tombol/detik)
0,75
22,545408 ms (186,038058 yuta tombol/detik)
0,78
26,053312 ms (160,989275 yuta tombol/detik)
0,81
31,895008 ms (131,503463 yuta tombol/detik)
0,84
42,103294 ms (99,619378 yuta tombol/detik)
0,87
61,849056 ms (67,815164 yuta tombol/detik)
0,90
105,695999 ms (39,682713 yuta tombol/detik)
0,94
240,204636 ms (17,461378 yuta tombol/detik)
Nalika panggunaan mundhak, kinerja mudhun. Iki ora dikarepake ing akeh kasus. Yen aplikasi nglebokake unsur menyang meja banjur dibuwang (contone, nalika ngetung tembung ing buku), mula iki ora dadi masalah. Nanging yen aplikasi nggunakake tabel hash sing wis suwe (contone, ing editor grafis kanggo nyimpen bagean gambar sing ora kosong ing ngendi pangguna asring nyisipake lan mbusak informasi), mula prilaku iki bisa dadi masalah.
Lan ngukur ambane probing tabel hash sawise 64 yuta sisipan (faktor pemanfaatan 0,5). Rata-rata ambane ana 0,4774, supaya paling tombol padha salah siji ing slot paling apik utawa siji slot adoh saka posisi paling apik. Kedalaman swara maksimal yaiku 60.
Aku banjur ngukur ambane probing ing meja karo 124 yuta sisipan (faktor pemanfaatan 0,97). Rata-rata ambane wis 10,1757, lan maksimum - 6474 (!!). Kinerja sensing linear mudhun banget ing tingkat panggunaan sing dhuwur.
Sing paling apik kanggo njaga tingkat panggunaan tabel hash iki sithik. Nanging banjur kita nambah kinerja ing beyo saka konsumsi memori. Untunge, ing kasus tombol lan nilai 32-bit, iki bisa dibenerake. Yen ing conto ing ndhuwur, ing meja karo 128 yuta unsur, kita tetep faktor pemanfaatan 0,25, banjur kita bisa manggonake ora luwih saka 32 yuta unsur ing, lan isih 96 yuta slot bakal ilang - 8 bita kanggo saben pasangan. , 768 MB ilang memori.
Wigati dimangerteni manawa kita ngomong babagan mundhut memori kertu video, sing minangka sumber daya sing luwih larang tinimbang memori sistem. Sanajan umume kertu grafis desktop modern sing ndhukung CUDA duwe paling ora 4 GB memori (nalika nulis, NVIDIA 2080 Ti duwe 11 GB), mesthine ora dadi keputusan sing paling wicaksana kanggo ngilangi jumlah kasebut.
Mengko aku bakal nulis liyane babagan nggawe tabel hash kanggo kertu video sing ora duwe masalah karo ambane probing, uga cara kanggo nggunakake maneh slot dibusak.
Pangukuran ambane swara
Kanggo nemtokake kedalaman probing kunci, kita bisa ngekstrak hash kunci (indeks tabel ideal) saka indeks tabel sing nyata:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Amarga saka Piandel loro loro kang nglengkapi nomer binar lan kasunyatan sing kapasitas Tabel hash loro kanggo daya loro, pendekatan iki bakal bisa malah nalika indeks tombol dipindhah menyang awal Tabel. Ayo dadi njupuk tombol sing hash kanggo 1, nanging dilebokake menyang slot 3. Banjur kanggo meja karo kapasitas 4 kita njaluk (3 — 1) & 3, sing padha karo 2.
kesimpulan
Yen sampeyan duwe pitakonan utawa komentar, please email kula ing utawa mbukak topik anyar ing.
Kode iki ditulis kanthi inspirasi saka artikel sing apik banget:
Ing mangsa ngarep, aku bakal terus nulis babagan implementasi tabel hash kanggo kertu video lan nganalisa kinerjae. Rencanaku kalebu chaining, Robin Hood hashing, lan cuckoo hashing nggunakake operasi atom ing struktur data sing ramah GPU.
Source: www.habr.com