Men uni Github-ga joylashtirdim .
Bu soniyada yuz millionlab qo'shimchalarni qayta ishlashga qodir oddiy GPU xesh jadvali. Mening NVIDIA GTX 1060 noutbukimda kod 64 milodiy tezlikda 210 million tasodifiy yaratilgan kalit-qiymat juftlarini kiritadi va taxminan 32 milodiyda 64 million juftni olib tashlaydi.
Ya'ni, noutbukda tezlik taxminan 300 million insert / sek va 500 million o'chirish / sek.
Jadval CUDA da yozilgan, garchi xuddi shu texnika HLSL yoki GLSL uchun qo'llanilishi mumkin. Amalga oshirish video kartada yuqori ishlashni ta'minlash uchun bir nechta cheklovlarga ega:
- Faqat 32-bitli kalitlar va bir xil qiymatlar qayta ishlanadi.
- Xesh jadvali qat'iy belgilangan o'lchamga ega.
- Va bu o'lcham kuchga ikkiga teng bo'lishi kerak.
Kalitlar va qiymatlar uchun siz oddiy ajratuvchi markerni zaxiralashingiz kerak (yuqoridagi kodda bu 0xffffffff).
Qulflarsiz xesh jadvali
Xesh-jadval bilan ochiq manzillashdan foydalanadi , ya'ni bu oddiygina xotirada saqlanadigan va yuqori kesh ishlashiga ega bo'lgan kalit-qiymat juftliklari massividir. Bog'langan ro'yxatdagi ko'rsatgichni qidirishni o'z ichiga olgan zanjirlash uchun ham xuddi shunday deb bo'lmaydi. Xesh-jadval elementlarni saqlaydigan oddiy massivdir KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Jadvalning o'lchami tub son emas, ikkining kuchidir, chunki pow2/VA niqobini qo'llash uchun bitta tezkor ko'rsatma yetarli, lekin modul operatori ancha sekinroq. Bu chiziqli zondlashda muhim ahamiyatga ega, chunki chiziqli jadvalni qidirishda slot indeksi har bir uyaga o'ralgan bo'lishi kerak. Va natijada, operatsiya narxi har bir uyaga modul qo'shiladi.
Jadval har bir element uchun kalit va qiymatni saqlaydi, kalitning xeshini emas. Jadvalda faqat 32 bitli kalitlar saqlanganligi sababli, xesh juda tez hisoblab chiqiladi. Yuqoridagi kod Murmur3 xesh-dan foydalanadi, u faqat bir nechta siljishlarni, XORlarni va ko'paytirishni amalga oshiradi.
Xesh jadvali xotira tartibiga bog'liq bo'lmagan qulflashni himoya qilish usullaridan foydalanadi. Ba'zi yozish operatsiyalari boshqa bunday operatsiyalarning tartibini buzsa ham, xesh jadvali hali ham to'g'ri holatni saqlab qoladi. Bu haqda quyida gaplashamiz. Texnika bir vaqtning o'zida minglab iplarni boshqaradigan video kartalar bilan ajoyib ishlaydi.
Xesh jadvalidagi kalitlar va qiymatlar bo'sh holga keltiriladi.
Kod 64 bitli kalitlar va qiymatlarni boshqarish uchun o'zgartirilishi mumkin. Kalitlar atomik o'qish, yozish va taqqoslash va almashtirish operatsiyalarini talab qiladi. Va qiymatlar atomik o'qish va yozish operatsiyalarini talab qiladi. Yaxshiyamki, CUDA-da 32 va 64 bitli qiymatlar uchun o'qish-yozish operatsiyalari tabiiy ravishda tekislangan bo'lsa, atomdir (pastga qarang). ) va zamonaviy video kartalar 64-bitli atom taqqoslash va almashish operatsiyalarini qo'llab-quvvatlaydi. Albatta, 64 bitga o'tishda ishlash biroz pasayadi.
Xesh-jadval holati
Xesh-jadvaldagi har bir kalit-qiymat juftligi to'rtta holatdan biriga ega bo'lishi mumkin:
- Kalit va qiymat bo'sh. Bu holatda, xesh jadvali ishga tushiriladi.
- Kalit yozildi, lekin qiymati hali yozilmagan. Agar boshqa mavzu hozirda ma'lumotlarni o'qiyotgan bo'lsa, u bo'sh qaytadi. Bu normal holat, agar boshqa ijro chizig'i biroz oldin ishlagan bo'lsa, xuddi shu narsa sodir bo'lardi va biz bir vaqtning o'zida ma'lumotlar tuzilishi haqida gapiramiz.
- Kalit ham, qiymat ham yozib olinadi.
- Qiymat boshqa bajarilish oqimlari uchun mavjud, ammo kalit hali yo'q. Bu CUDA dasturlash modeli bo'sh tartiblangan xotira modeliga ega bo'lgani uchun sodir bo'lishi mumkin. Bu normal holat; har qanday holatda ham, qiymat endi bo'lmasa ham, kalit hali ham bo'sh.
Muhim nuance shundaki, kalit uyaga yozilgandan so'ng, u endi harakat qilmaydi - hatto kalit o'chirilgan bo'lsa ham, biz bu haqda quyida gaplashamiz.
Xesh-jadval kodi hatto xotirani o'qish va yozish tartibi noma'lum bo'lgan erkin tartiblangan xotira modellari bilan ishlaydi. Xesh-jadvalga qo'shish, qidirish va o'chirishni ko'rib chiqsak, har bir kalit-qiymat juftligi yuqorida tavsiflangan to'rtta holatdan birida ekanligini unutmang.
Xesh jadvaliga kiritish
Xesh-jadvalga kalit-qiymat juftlarini kiritadigan CUDA funktsiyasi quyidagicha ko'rinadi:
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);
}
}
Kalitni kiritish uchun kod kiritilgan kalitning xeshidan boshlab xesh jadvali qatorida takrorlanadi. Massivdagi har bir slot atomik taqqoslash va almashtirish operatsiyasini bajaradi, bu esa ushbu slotdagi kalitni bo'sh bilan solishtiradi. Agar nomuvofiqlik aniqlansa, uyadagi kalit kiritilgan kalit bilan yangilanadi va keyin asl slot kaliti qaytariladi. Agar bu asl kalit bo'sh bo'lsa yoki kiritilgan kalitga to'g'ri kelsa, u holda kod kiritish uchun mos bo'sh joy topdi va kiritilgan qiymatni uyaga kiritdi.
Agar bitta yadro chaqiruvida bo'lsa gpu_hashtable_insert() bir xil kalitga ega bo'lgan bir nechta elementlar mavjud, keyin ularning har qanday qiymatlari kalit uyasiga yozilishi mumkin. Bu odatiy hisoblanadi: qo'ng'iroq paytida kalit-qiymat yozishlaridan biri muvaffaqiyatli bo'ladi, ammo bularning barchasi bir nechta bajarilish bosqichlarida parallel ravishda sodir bo'lganligi sababli, biz qaysi xotira yozish oxirgi bo'lishini oldindan aytib bo'lmaydi.
Xesh jadvalini qidirish
Kalitlarni qidirish uchun kod:
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);
}
}
Jadvalda saqlangan kalitning qiymatini topish uchun biz izlayotgan kalitning xeshidan boshlab massivni takrorlaymiz. Har bir slotda kalit biz izlayotgan kalit ekanligini tekshiramiz va agar shunday bo'lsa, biz uning qiymatini qaytaramiz. Shuningdek, biz kalit bo'sh yoki yo'qligini tekshiramiz va agar shunday bo'lsa, qidiruvni to'xtatamiz.
Agar kalitni topa olmasak, kod bo'sh qiymatni qaytaradi.
Ushbu qidiruv operatsiyalarining barchasi qo'shish va o'chirish orqali bir vaqtning o'zida amalga oshirilishi mumkin. Jadvaldagi har bir juftlik oqim uchun yuqorida tavsiflangan to'rtta holatdan biriga ega bo'ladi.
Xesh-jadvalda o'chirish
Kalitlarni o'chirish uchun kod:
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);
}
}
Kalitni o'chirish noodatiy tarzda amalga oshiriladi: biz kalitni jadvalda qoldiramiz va uning qiymatini (kalitning o'zi emas) bo'sh deb belgilaymiz. Ushbu kod juda o'xshash lookup(), bundan mustasno, agar kalitda moslik topilsa, uning qiymatini bo'sh qiladi.
Yuqorida aytib o'tilganidek, kalit uyaga yozilgach, u boshqa joyga ko'chirilmaydi. Jadvaldan element o'chirilganda ham, kalit joyida qoladi, uning qiymati shunchaki bo'sh bo'ladi. Bu shuni anglatadiki, biz slot qiymati uchun atomik yozish operatsiyasidan foydalanishimiz shart emas, chunki joriy qiymat bo'sh yoki bo'lmasligi muhim emas - u hali ham bo'sh bo'lib qoladi.
Xesh-jadval hajmini o'zgartirish
Kattaroq jadval yaratish va unga eski jadvaldan boʻsh boʻlmagan elementlarni kiritish orqali xesh-jadval hajmini oʻzgartirishingiz mumkin. Men bu funksiyani amalga oshirmadim, chunki namuna kodini oddiy saqlashni xohladim. Bundan tashqari, CUDA dasturlarida xotirani ajratish ko'pincha CUDA yadrosida emas, balki xost kodida amalga oshiriladi.
maqola bunday qulfdan himoyalangan ma'lumotlar strukturasini qanday o'zgartirishni tasvirlaydi.
Raqobatbardoshlik
Yuqoridagi funksiya kod qismlarida gpu_hashtable_insert(), _lookup() и _delete() bir vaqtning o'zida bitta kalit-qiymat juftligini qayta ishlash. Va pastroq gpu_hashtable_insert(), _lookup() и _delete() parallel ravishda bir qator juftliklarni qayta ishlash, har bir juftlik alohida GPU ijro ipida:
// 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);
}
}
Qulfga chidamli xesh jadvali bir vaqtda qo'shish, qidirish va o'chirishni qo'llab-quvvatlaydi. Kalit-qiymat juftligi har doim to'rtta holatdan birida bo'lgani va kalitlar harakat qilmasligi sababli, jadval turli xil operatsiyalar bir vaqtning o'zida ishlatilganda ham to'g'riligini kafolatlaydi.
Biroq, agar biz kiritishlar va o'chirishlar to'plamini parallel ravishda qayta ishlasak va agar juftliklarning kirish massivida takroriy kalitlar bo'lsa, biz qaysi juftliklar "yutishini" taxmin qila olmaymiz - xesh jadvaliga oxirgi yoziladi. Aytaylik, biz kiritish kodini juftlik kiritish massivi bilan chaqirdik A/0 B/1 A/2 C/3 A/4. Kod tugagach, juftlashadi B/1 и C/3 jadvalda mavjud bo'lishi kafolatlanadi, lekin bir vaqtning o'zida har qanday juftlik unda paydo bo'ladi A/0, A/2 yoki A/4. Bu muammo bo'lishi mumkin yoki bo'lmasligi mumkin - barchasi dasturga bog'liq. Kirish massivida takroriy kalitlar yo'qligini oldindan bilishingiz mumkin yoki qaysi qiymat oxirgi yozilganiga ahamiyat bermasligingiz mumkin.
Agar bu siz uchun muammo bo'lsa, unda siz ikki nusxadagi juftlarni turli CUDA tizim chaqiruvlariga ajratishingiz kerak. CUDA da yadroni chaqiruvchi har qanday operatsiya har doim navbatdagi yadro chaqiruvidan oldin tugallanadi (hech bo'lmaganda bitta ip ichida. Turli xil oqimlarda yadrolar parallel ravishda bajariladi). Yuqoridagi misolda, agar siz bitta yadro bilan qo'ng'iroq qilsangiz A/0 B/1 A/2 C/3, va ikkinchisi bilan A/4, keyin kalit A qiymatini oladi 4.
Endi funksiyalar kerakmi yoki yo'qmi haqida gapiraylik lookup() и delete() xesh jadvalidagi juftliklar qatoriga oddiy yoki o'zgaruvchan ko'rsatgichdan foydalaning. Quyidagilarni bildiradi:
Kompilyator global yoki umumiy xotiraga o'qish va yozishni optimallashtirishni tanlashi mumkin... Ushbu optimallashtirishlarni kalit so'z yordamida o'chirib qo'yish mumkin.
volatile: ... bu o'zgaruvchiga har qanday havola haqiqiy xotira o'qish yoki yozish ko'rsatmasi ichiga kompilyatsiya qilinadi.
To'g'riligi haqidagi fikrlar qo'llashni talab qilmaydi volatile. Agar ijro chizig'i oldingi o'qish operatsiyasidan keshlangan qiymatdan foydalansa, u biroz eskirgan ma'lumotlardan foydalanadi. Ammo shunga qaramay, bu yadro chaqiruvining ma'lum bir daqiqasida xesh jadvalining to'g'ri holatidan olingan ma'lumot. Agar siz eng so'nggi ma'lumotlardan foydalanishingiz kerak bo'lsa, indeksdan foydalanishingiz mumkin volatile, lekin keyin ishlash biroz pasayadi: mening testlarimga ko'ra, 32 million elementni o'chirishda tezlik 500 million o'chirish / sek dan 450 million o'chirish / sek gacha kamaydi.
unumdorlik
Sinovda 64 million elementni kiritish va ulardan 32 millionini o'chirish uchun raqobat mavjud std::unordered_map va GPU uchun xesh jadvali deyarli yo'q:
std::unordered_map elementlarni kiritish va olib tashlash va keyin ularni bo'shatish uchun 70 ms sarfladi unordered_map (millionlab elementlardan xalos bo'lish juda ko'p vaqtni oladi, chunki ichkarida unordered_map bir nechta xotira ajratish amalga oshiriladi). Rostini aytsam, std:unordered_map butunlay boshqacha cheklovlar. Bu bitta protsessorning bajarilishi, har qanday o'lchamdagi kalit-qiymatlarni qo'llab-quvvatlaydi, yuqori foydalanish tezligida yaxshi ishlaydi va bir nechta o'chirishdan keyin barqaror ishlashni ko'rsatadi.
GPU va dasturlararo aloqa uchun xesh-jadvalning davomiyligi 984 milodiyni tashkil etdi. Bunga jadvalni xotiraga joylashtirish va uni oʻchirish (1 Gb xotirani bir marta ajratish, bu CUDAda biroz vaqt talab etadi), elementlarni kiritish va oʻchirish hamda ular ustida takrorlash vaqtini oʻz ichiga oladi. Video karta xotirasiga va undan olingan barcha nusxalar ham hisobga olinadi.
Xesh-jadvalning o'zi bajarish uchun 271 milodiy vaqtni oldi. Bu video kartaning elementlarni kiritish va o'chirish uchun sarflagan vaqtini o'z ichiga oladi va xotiraga nusxa ko'chirish va natijada olingan jadvalni takrorlash uchun sarflangan vaqtni hisobga olmaydi. Agar GPU jadvali uzoq vaqt yashasa yoki xesh-jadval to'liq video karta xotirasida bo'lsa (masalan, markaziy protsessor emas, balki boshqa GPU kodi tomonidan ishlatiladigan xesh jadvalini yaratish uchun), keyin test natijasi tegishli.
Video karta uchun xesh jadvali yuqori o'tkazuvchanlik va faol parallelizatsiya tufayli yuqori ishlashni namoyish etadi.
kamchiliklar
Xesh-jadval arxitekturasida bilish kerak bo'lgan bir nechta muammolar mavjud:
- Chiziqli probing klasterlash bilan to'sqinlik qiladi, bu esa jadvaldagi kalitlarni mukammal darajadan kamroq joylashtirishga olib keladi.
- Funktsiya yordamida kalitlar olib tashlanmaydi
deleteva vaqt o'tishi bilan ular stolni chalkashtirib yuborishadi.
Natijada, xesh-jadvalning ishlashi asta-sekin yomonlashishi mumkin, ayniqsa u uzoq vaqt davomida mavjud bo'lsa va ko'plab qo'shimchalar va o'chirishlar mavjud bo'lsa. Ushbu kamchiliklarni yumshatishning bir usuli - foydalanish darajasi juda past bo'lgan yangi jadvalga o'tish va qayta tiklash paytida olib tashlangan kalitlarni filtrlash.
Ta'riflangan muammolarni tasvirlash uchun men 128 million elementdan iborat jadval yaratish uchun yuqoridagi koddan foydalanaman va 4 million slotni to'ldirgunimcha 124 million elementni aylantiraman (foydalanish darajasi taxminan 0,96). Mana natijalar jadvali, har bir satr bitta xesh jadvaliga 4 million yangi elementni kiritish uchun CUDA yadro chaqiruvidir:
Foydalanish darajasi
Qo‘shish davomiyligi 4 194 304 element
0,00
11,608448 ms (361,314798 million kalit/sek.)
0,03
11,751424 ms (356,918799 million kalit/sek.)
0,06
11,942592 ms (351,205515 million kalit/sek.)
0,09
12,081120 ms (347,178429 million kalit/sek.)
0,12
12,242560 ms (342,600233 million kalit/sek.)
0,16
12,396448 ms (338,347235 million kalit/sek.)
0,19
12,533024 ms (334,660176 million kalit/sek.)
0,22
12,703328 ms (330,173626 million kalit/sek.)
0,25
12,884512 ms (325,530693 million kalit/sek.)
0,28
13,033472 ms (321,810182 million kalit/sek.)
0,31
13,239296 ms (316,807174 million kalit/sek.)
0,34
13,392448 ms (313,184256 million kalit/sek.)
0,37
13,624000 ms (307,861434 million kalit/sek.)
0,41
13,875520 ms (302,280855 million kalit/sek.)
0,44
14,126528 ms (296,909756 million kalit/sek.)
0,47
14,399328 ms (291,284699 million kalit/sek.)
0,50
14,690304 ms (285,515123 million kalit/sek.)
0,53
15,039136 ms (278,892623 million kalit/sek.)
0,56
15,478656 ms (270,973402 million kalit/sek.)
0,59
15,985664 ms (262,379092 million kalit/sek.)
0,62
16,668673 ms (251,627968 million kalit/sek.)
0,66
17,587200 ms (238,486174 million kalit/sek.)
0,69
18,690048 ms (224,413765 million kalit/sek.)
0,72
20,278816 ms (206,831789 million kalit/sek.)
0,75
22,545408 ms (186,038058 million kalit/sek.)
0,78
26,053312 ms (160,989275 million kalit/sek.)
0,81
31,895008 ms (131,503463 million kalit/sek.)
0,84
42,103294 ms (99,619378 million kalit/sek.)
0,87
61,849056 ms (67,815164 million kalit/sek.)
0,90
105,695999 ms (39,682713 million kalit/sek.)
0,94
240,204636 ms (17,461378 million kalit/sek.)
Foydalanish ortishi bilan unumdorlik pasayadi. Aksariyat hollarda bu istalmagan. Agar dastur jadvalga elementlarni kiritsa va keyin ularni tashlab qo'ysa (masalan, kitobdagi so'zlarni sanashda), unda bu muammo emas. Ammo agar ilova uzoq muddatli xesh-jadvaldan foydalansa (masalan, foydalanuvchi ma'lumotni tez-tez kiritadigan va o'chiradigan tasvirlarning bo'sh bo'lmagan qismlarini saqlash uchun grafik muharrirda), unda bu xatti-harakatlar muammoli bo'lishi mumkin.
Va 64 million qo'shimchadan keyin xesh jadvalini tekshirish chuqurligini o'lchadi (foydalanish koeffitsienti 0,5). O'rtacha chuqurlik 0,4774 ni tashkil etdi, shuning uchun ko'pchilik kalitlar mumkin bo'lgan eng yaxshi uyada yoki eng yaxshi holatdan bir uyasi uzoqda edi. Maksimal tovush chuqurligi 60 edi.
Keyin men 124 million qo'shimchali stolda zondlash chuqurligini o'lchadim (foydalanish koeffitsienti 0,97). O'rtacha chuqurlik allaqachon 10,1757, maksimal - 6474 (!!). Yuqori foydalanish stavkalarida chiziqli sezish samaradorligi sezilarli darajada pasayadi.
Ushbu xesh-jadvaldan foydalanish darajasini past darajada ushlab turish yaxshidir. Ammo keyin biz xotira iste'moli hisobiga ishlashni oshiramiz. Yaxshiyamki, 32-bitli kalitlar va qiymatlar bo'lsa, buni oqlash mumkin. Agar yuqoridagi misolda 128 million elementdan iborat jadvalda biz foydalanish koeffitsientini 0,25 ga teng ushlab tursak, unda biz 32 milliondan ortiq elementni joylashtira olmaymiz va qolgan 96 million slot yo'qoladi - har bir juftlik uchun 8 bayt. , 768 MB yo'qolgan xotira.
E'tibor bering, biz tizim xotirasidan ko'ra qimmatroq manba bo'lgan video karta xotirasining yo'qolishi haqida gapiramiz. CUDA-ni qo'llab-quvvatlaydigan zamonaviy ish stoli grafik kartalarining ko'pchiligi kamida 4 Gb xotiraga ega bo'lsa-da (yozilish vaqtida NVIDIA 2080 Ti 11 Gb xotiraga ega), bu kabi miqdorlarni yo'qotish hali ham eng oqilona qaror bo'lmaydi.
Keyinchalik men problash chuqurligi bilan bog'liq muammolar bo'lmagan video kartalar uchun xesh jadvallarini yaratish, shuningdek, o'chirilgan slotlarni qayta ishlatish usullari haqida ko'proq yozaman.
Ovoz chuqurligini o'lchash
Kalitni tekshirish chuqurligini aniqlash uchun biz uning haqiqiy jadval indeksidan kalit xeshini (uning ideal jadval indeksini) ajratib olishimiz mumkin:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Ikkining ikkitasini to'ldiruvchi ikkilik sonlarning sehrliligi va xesh-jadvalning sig'imi ikkitaning ikkita darajasiga teng bo'lganligi sababli, bu yondashuv kalit indeks jadvalning boshiga ko'chirilganda ham ishlaydi. Keling, 1 ga xeshlangan, lekin 3-uyaga kiritilgan kalitni olaylik. Keyin 4 sig'imli stol uchun biz olamiz (3 — 1) & 3, bu 2 ga teng.
xulosa
Savollaringiz yoki sharhlaringiz bo'lsa, iltimos, menga elektron pochta orqali yuboring yoki yangi mavzu oching.
Ushbu kod ajoyib maqolalardan ilhomlanib yozilgan:
Kelajakda men video kartalar uchun hash-jadvalni amalga oshirish haqida yozishni davom ettiraman va ularning ishlashini tahlil qilaman. Mening rejalarimga zanjirband qilish, Robin Gud xeshing va GPU bilan mos keladigan ma'lumotlar tuzilmalarida atom operatsiyalaridan foydalangan holda kuku xeshing kiradi.
Manba: www.habr.com