Mən bunu Github-da yerləşdirdim .
Bu, saniyədə yüz milyonlarla əlavəni emal edə bilən sadə GPU hash cədvəlidir. NVIDIA GTX 1060 noutbukumda kod təxminən 64 ms-də təsadüfi yaradılan 210 milyon açar-dəyər cütünü daxil edir və təxminən 32 ms-də 64 milyon cütü silir.
Yəni, noutbukda sürət təxminən 300 milyon insert/san və 500 milyon silmə/san təşkil edir.
Cədvəl CUDA-da yazılmışdır, baxmayaraq ki, eyni texnika HLSL və ya GLSL-də tətbiq oluna bilər. Tətbiq video kartda yüksək performansı təmin etmək üçün bir sıra məhdudiyyətlərə malikdir:
- Yalnız 32 bitlik açarlar və eyni dəyərlər işlənir.
- Hash cədvəlinin sabit ölçüsü var.
- Və bu ölçü gücə ikiyə bərabər olmalıdır.
Düymələr və dəyərlər üçün sadə ayırıcı marker rezerv etməlisiniz (yuxarıdakı kodda bu 0xffffffff-dir).
Kilidləri olmayan hash cədvəli
Hash cədvəli ilə açıq ünvanlama istifadə olunur , yəni yaddaşda saxlanılan və üstün keş performansına malik olan sadəcə açar-dəyər cütləri massividir. Eyni şeyi əlaqəli siyahıda göstərici axtarmağı nəzərdə tutan zəncirləmə üçün demək olmaz. Hash cədvəli elementləri saxlayan sadə massivdir KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Cədvəlin ölçüsü ikinin gücüdür, sadə rəqəm deyil, çünki pow2/AND maskasını tətbiq etmək üçün bir sürətli təlimat kifayətdir, lakin modul operatoru daha yavaşdır. Bu, xətti zondlama zamanı vacibdir, çünki xətti cədvəl axtarışında yuva indeksi hər bir yuvaya bükülməlidir. Və nəticədə əməliyyatın dəyəri hər bir yuvaya modul əlavə olunur.
Cədvəl yalnız açarı və hər bir element üçün dəyəri saxlayır, açarın hashını deyil. Cədvəl yalnız 32 bitlik açarları saxladığından, hash çox tez hesablanır. Yuxarıdakı kod yalnız bir neçə növbə, XOR və vurma yerinə yetirən Murmur3 hashından istifadə edir.
Hash cədvəli yaddaş sırasından asılı olmayan kilidləmə mühafizə üsullarından istifadə edir. Bəzi yazma əməliyyatları digər bu cür əməliyyatların sırasını pozsa belə, hash cədvəli hələ də düzgün vəziyyəti saxlayacaqdır. Bu barədə aşağıda danışacağıq. Texnika minlərlə ipi eyni vaxtda işlədən video kartlarla əla işləyir.
Hash cədvəlindəki açarlar və dəyərlər boş yerə işə salınır.
Kod 64 bitlik açarları və dəyərləri idarə etmək üçün dəyişdirilə bilər. Açarlar atom oxu, yazma və müqayisə və dəyişdirmə əməliyyatlarını tələb edir. Və dəyərlər atom oxuma və yazma əməliyyatlarını tələb edir. Xoşbəxtlikdən, CUDA-da 32 və 64 bitlik dəyərlər üçün oxuma-yazma əməliyyatları təbii olaraq hizalandığı müddətcə atomikdir (aşağıya bax). ) və müasir video kartlar 64 bitlik atom müqayisə və mübadilə əməliyyatlarını dəstəkləyir. Təbii ki, 64 bitə keçərkən performans bir qədər azalacaq.
Hash cədvəlinin vəziyyəti
Hash cədvəlindəki hər bir açar-dəyər cütü dörd vəziyyətdən birinə malik ola bilər:
- Açar və dəyər boşdur. Bu vəziyyətdə, hash cədvəli işə salınır.
- Açar yazılıb, amma dəyəri hələ yazılmayıb. Əgər başqa mövzu hazırda məlumatları oxuyursa, o, boş qaytarır. Bu normaldır, başqa bir icra xətti bir az əvvəl işləsəydi, eyni şey baş verərdi və biz paralel məlumat strukturundan danışırıq.
- Həm açar, həm də dəyər qeyd olunur.
- Dəyər icranın digər mövzuları üçün əlçatandır, lakin açar hələ yoxdur. Bu, CUDA proqramlaşdırma modelinin zəif sifariş edilmiş yaddaş modelinə malik olması səbəbindən baş verə bilər. Bu normaldır; istənilən halda, dəyər artıq olmasa belə, açar hələ də boşdur.
Əhəmiyyətli bir nüans odur ki, açar yuvaya yazıldıqdan sonra o, artıq hərəkət etmir - açar silinsə belə, aşağıda bu barədə danışacağıq.
Hash-cədvəl kodu hətta yaddaşın oxunma və yazılma ardıcıllığının bilinmədiyi boş sifarişli yaddaş modelləri ilə də işləyir. Hash cədvəlində daxiletmə, axtarış və silinmə ilə bağlı məsələlərə baxarkən yadda saxlayın ki, hər bir açar-dəyər cütü yuxarıda təsvir edilən dörd vəziyyətdən birindədir.
Hash cədvəlinə daxil edilir
Açar-dəyər cütlərini hash cədvəlinə daxil edən CUDA funksiyası belə görünür:
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);
}
}
Açarı daxil etmək üçün kod daxil edilmiş açarın heşindən başlayaraq hash cədvəli massivini təkrarlayır. Massivdəki hər bir yuva atomik müqayisə və dəyişdirmə əməliyyatını yerinə yetirir ki, bu da həmin yuvadakı açarı boş yerə müqayisə edir. Uyğunsuzluq aşkar edilərsə, yuvadakı açar daxil edilmiş açarla yenilənir və sonra orijinal yuva açarı qaytarılır. Bu orijinal açar boş idisə və ya daxil edilmiş açarla uyğun gəlirsə, kod daxil etmək üçün uyğun yuva tapdı və daxil edilmiş dəyəri yuvaya daxil etdi.
Bir kernel çağırışı varsa gpu_hashtable_insert() eyni açarı olan bir neçə element var, onda onların hər hansı bir dəyəri açar yuvasına yazıla bilər. Bu normal hesab olunur: zəng zamanı açar-dəyər yazılarından biri uğur qazanacaq, lakin bütün bunlar bir neçə icra xətti daxilində paralel olaraq baş verdiyindən, hansı yaddaş yazısının sonuncu olacağını təxmin edə bilmərik.
Hash cədvəlinin axtarışı
Axtarış açarları üçün 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);
}
}
Cədvəldə saxlanılan açarın dəyərini tapmaq üçün biz axtardığımız açarın hashından başlayaraq massiv üzrə təkrar edirik. Hər bir yuvada açarın axtardığımız açar olub-olmadığını yoxlayırıq və əgər belədirsə, onun dəyərini qaytarırıq. Açarın boş olub olmadığını da yoxlayırıq və əgər belədirsə, axtarışı dayandırırıq.
Əgər açarı tapa bilməsək, kod boş bir dəyər qaytarır.
Bütün bu axtarış əməliyyatları əlavələr və silinmələr vasitəsilə eyni vaxtda həyata keçirilə bilər. Cədvəldəki hər bir cüt axın üçün yuxarıda təsvir edilən dörd vəziyyətdən birinə malik olacaq.
Hash cədvəlində silinir
Açarları silmək üçün 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);
}
}
Açarın silinməsi qeyri-adi üsulla həyata keçirilir: biz açarı cədvəldə qoyub onun dəyərini (açarın özünü deyil) boş kimi qeyd edirik. Bu kod çox oxşardır lookup(), istisna olmaqla, açarda uyğunluq aşkar edildikdə, onun dəyərini boş edir.
Yuxarıda qeyd edildiyi kimi, bir açar bir yuvaya yazıldıqdan sonra o, artıq köçürülmür. Cədvəldən element silindikdə belə açar yerində qalır, onun dəyəri sadəcə boş olur. Bu o deməkdir ki, yuva dəyəri üçün atom yazma əməliyyatından istifadə etməyə ehtiyac yoxdur, çünki cari dəyərin boş olub-olmamasının əhəmiyyəti yoxdur - o, yenə də boşalacaq.
Hash cədvəlinin ölçüsünü dəyişdirin
Daha böyük bir cədvəl yaradaraq və ona köhnə cədvəldən boş olmayan elementlər daxil etməklə hash cədvəlinin ölçüsünü dəyişə bilərsiniz. Mən bu funksiyanı tətbiq etmədim, çünki nümunə kodunu sadə saxlamaq istəyirdim. Üstəlik, CUDA proqramlarında yaddaşın ayrılması çox vaxt CUDA nüvəsində deyil, host kodunda həyata keçirilir.
Məqalədə belə bir kilidlə qorunan məlumat strukturunun necə dəyişdiriləcəyini təsvir edir.
Rəqabət qabiliyyəti
Yuxarıdakı funksiya kodu parçalarında gpu_hashtable_insert(), _lookup() и _delete() bir anda bir açar-dəyər cütünü emal edin. Və aşağı gpu_hashtable_insert(), _lookup() и _delete() paralel olaraq bir sıra cütləri emal edin, hər bir cüt ayrı GPU icra ipində:
// 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);
}
}
Kilidə davamlı hash cədvəli paralel daxiletmələri, axtarışları və silmələri dəstəkləyir. Açar-dəyər cütləri həmişə dörd vəziyyətdən birində olduğundan və düymələr yerindən tərpənmədiyindən, müxtəlif növ əməliyyatlar eyni vaxtda istifadə edildikdə belə cədvəl düzgünlüyünə zəmanət verir.
Bununla belə, əgər biz əlavələr və silmələr dəstini paralel olaraq emal etsək və cütlərin giriş massivində dublikat açarlar varsa, o zaman hansı cütlərin “qalib” olacağını təxmin edə bilməyəcəyik — hash cədvəlinə sonuncu yazılacaq. Deyək ki, biz daxiletmə kodunu cütlərin giriş massivi ilə çağırdıq A/0 B/1 A/2 C/3 A/4. Kod tamamlandıqda, cütləşir B/1 и C/3 cədvəldə mövcud olacağına zəmanət verilir, lakin eyni zamanda cütlərdən hər hansı biri onda görünəcəkdir A/0, A/2 və ya A/4. Bu problem ola bilər və ya olmaya bilər - hamısı tətbiqdən asılıdır. Siz əvvəlcədən bilə bilərsiniz ki, giriş massivində dublikat düymələr yoxdur və ya hansı dəyərin sonuncu yazılmasına əhəmiyyət verməyə bilərsiniz.
Əgər bu sizin üçün problemdirsə, onda siz dublikat cütlərini müxtəlif CUDA sistem zənglərinə ayırmalısınız. CUDA-da nüvəni çağıran istənilən əməliyyat həmişə növbəti kernel çağırışından əvvəl tamamlanır (ən azı bir iplik daxilində. Müxtəlif ipliklərdə ləpələr paralel olaraq yerinə yetirilir). Yuxarıdakı misalda, bir kernel ilə zəng etsəniz A/0 B/1 A/2 C/3, digəri ilə A/4, sonra açar A dəyərini alacaq 4.
İndi funksiyaların lazım olub-olmaması barədə danışaq lookup() и delete() hash cədvəlindəki cütlər massivinə düz və ya dəyişkən göstəricidən istifadə edin. Bildirir ki:
Kompilyator qlobal və ya paylaşılan yaddaşa oxuma və yazmaları optimallaşdırmağı seçə bilər... Bu optimallaşdırmalar açar sözdən istifadə etməklə deaktiv edilə bilər.
volatile: ... bu dəyişənə hər hansı istinad real yaddaşın oxunması və ya yazılması təlimatında tərtib edilir.
Düzgünlük mülahizələri tətbiq tələb etmir volatile. Əgər icra başlığı əvvəlki oxunmuş əməliyyatdan keşlənmiş dəyərdən istifadə edirsə, o, bir qədər köhnəlmiş məlumatdan istifadə edəcək. Ancaq yenə də bu, nüvə çağırışının müəyyən bir anında hash cədvəlinin düzgün vəziyyətindən alınan məlumatdır. Ən son məlumatlardan istifadə etmək lazımdırsa, indeksdən istifadə edə bilərsiniz volatile, lakin sonra performans bir qədər azalacaq: testlərimə görə, 32 milyon elementi silərkən sürət 500 milyon silinmə/sandan 450 milyon silinmə/san-a qədər azaldı.
Məhsuldarlıq
Testdə 64 milyon elementin daxil edilməsi və onlardan 32 milyonunun silinməsi arasında rəqabət var std::unordered_map və GPU üçün praktiki olaraq heç bir hash cədvəli yoxdur:
std::unordered_map elementləri daxil etmək və çıxarmaq və sonra onları azad etmək üçün 70 ms sərf etmişdir unordered_map (milyonlarla elementdən qurtulmaq çox vaxt aparır, çünki içəridə unordered_map çox yaddaş ayrılması edilir). Düzünü desəm, std:unordered_map tamamilə fərqli məhdudiyyətlər. Bu, tək CPU-nun icra xəttidir, istənilən ölçülü açar-dəyərləri dəstəkləyir, yüksək istifadə dərəcələrində yaxşı işləyir və çoxsaylı silinmədən sonra sabit performans göstərir.
GPU və proqramlararası əlaqə üçün hash cədvəlinin müddəti 984 ms idi. Buraya cədvəlin yaddaşa yerləşdirilməsi və silinməsi (bir dəfə 1 GB yaddaşın ayrılması, bu, CUDA-da müəyyən vaxt tələb edir), elementlərin daxil edilməsi və silinməsi və onların üzərində təkrarlamaya sərf olunan vaxt daxildir. Video kartın yaddaşına daxil olan və ondan çıxan bütün nüsxələr də nəzərə alınır.
Hash cədvəlinin özü tamamlanması üçün 271 ms çəkdi. Buraya video kartın elementləri daxil etmək və silmək üçün sərf etdiyi vaxt daxildir və yaddaşa köçürmə və nəticədə ortaya çıxan cədvəl üzərində təkrarlama üçün sərf olunan vaxtı nəzərə almır. GPU cədvəli uzun müddət yaşayırsa və ya hash cədvəli tamamilə video kartın yaddaşındadırsa (məsələn, mərkəzi prosessor tərəfindən deyil, digər GPU kodu tərəfindən istifadə ediləcək bir hash cədvəli yaratmaq üçün), onda test nəticəsi müvafiqdir.
Video kart üçün hash cədvəli yüksək ötürmə qabiliyyəti və aktiv paralelləşdirmə sayəsində yüksək performans nümayiş etdirir.
Məhdudiyyətlər
Hash cədvəlinin arxitekturasında bir neçə məsələdən xəbərdar olmaq lazımdır:
- Xətti zondlama qruplaşma ilə maneə törədir, bu da cədvəldəki açarların mükəmməl olduğundan daha az yerləşdirilməsinə səbəb olur.
- Funksiyadan istifadə edərək düymələr çıxarılmır
deletevə zaman keçdikcə masanı darmadağın edirlər.
Nəticədə, hash cədvəlinin performansı, xüsusən də uzun müddət mövcud olduqda və çoxsaylı əlavələr və silmələr varsa, tədricən pisləşə bilər. Bu çatışmazlıqları azaltmağın bir yolu, kifayət qədər aşağı istifadə nisbəti ilə yeni bir cədvələ yenidən daxil etmək və rehashing zamanı çıxarılan düymələri süzgəcdən keçirməkdir.
Təsvir edilən məsələləri təsvir etmək üçün yuxarıdakı koddan istifadə edərək 128 milyon elementdən ibarət cədvəl yaradacağam və 4 milyon yuvanı doldurana qədər 124 milyon element arasında dövrə vuracağam (istifadə nisbəti təxminən 0,96). Budur nəticə cədvəli, hər bir sıra bir hash cədvəlinə 4 milyon yeni element daxil etmək üçün CUDA ləpə çağırışıdır:
İstifadə dərəcəsi
Daxiletmə müddəti 4 element
0,00
11,608448 ms (361,314798 milyon açar/san)
0,03
11,751424 ms (356,918799 milyon açar/san)
0,06
11,942592 ms (351,205515 milyon açar/san)
0,09
12,081120 ms (347,178429 milyon açar/san)
0,12
12,242560 ms (342,600233 milyon açar/san)
0,16
12,396448 ms (338,347235 milyon açar/san)
0,19
12,533024 ms (334,660176 milyon açar/san)
0,22
12,703328 ms (330,173626 milyon açar/san)
0,25
12,884512 ms (325,530693 milyon açar/san)
0,28
13,033472 ms (321,810182 milyon açar/san)
0,31
13,239296 ms (316,807174 milyon açar/san)
0,34
13,392448 ms (313,184256 milyon açar/san)
0,37
13,624000 ms (307,861434 milyon açar/san)
0,41
13,875520 ms (302,280855 milyon açar/san)
0,44
14,126528 ms (296,909756 milyon açar/san)
0,47
14,399328 ms (291,284699 milyon açar/san)
0,50
14,690304 ms (285,515123 milyon açar/san)
0,53
15,039136 ms (278,892623 milyon açar/san)
0,56
15,478656 ms (270,973402 milyon açar/san)
0,59
15,985664 ms (262,379092 milyon açar/san)
0,62
16,668673 ms (251,627968 milyon açar/san)
0,66
17,587200 ms (238,486174 milyon açar/san)
0,69
18,690048 ms (224,413765 milyon açar/san)
0,72
20,278816 ms (206,831789 milyon açar/san)
0,75
22,545408 ms (186,038058 milyon açar/san)
0,78
26,053312 ms (160,989275 milyon açar/san)
0,81
31,895008 ms (131,503463 milyon açar/san)
0,84
42,103294 ms (99,619378 milyon açar/san)
0,87
61,849056 ms (67,815164 milyon açar/san)
0,90
105,695999 ms (39,682713 milyon açar/san)
0,94
240,204636 ms (17,461378 milyon açar/san)
İstifadə artdıqca performans azalır. Bu, əksər hallarda arzuolunan deyil. Tətbiq cədvələ elementlər daxil edirsə və sonra onları ləğv edirsə (məsələn, kitabdakı sözləri sayarkən), onda bu problem deyil. Ancaq tətbiq uzunmüddətli hash cədvəlindən istifadə edirsə (məsələn, istifadəçinin məlumatı tez-tez daxil etdiyi və sildiyi şəkillərin boş olmayan hissələrini saxlamaq üçün qrafik redaktorda), onda bu davranış problemli ola bilər.
Və 64 milyon əlavədən sonra hash cədvəlinin zondlama dərinliyini ölçdü (istifadə əmsalı 0,5). Orta dərinlik 0,4774 idi, buna görə də əksər düymələr ya mümkün olan ən yaxşı yuvada, ya da ən yaxşı mövqedən bir yuvada idi. Maksimum səslənmə dərinliyi 60 idi.
Daha sonra 124 milyon əlavə ilə masada zondlama dərinliyini ölçdüm (istifadə əmsalı 0,97). Orta dərinlik artıq 10,1757, maksimum isə - 6474 (!!). Xətti algılama performansı yüksək istifadə dərəcələrində əhəmiyyətli dərəcədə azalır.
Bu hash cədvəlinin istifadə dərəcəsini aşağı saxlamaq yaxşıdır. Amma sonra yaddaş sərfiyyatı hesabına performansı artırırıq. Xoşbəxtlikdən, 32 bitlik açarlar və dəyərlər vəziyyətində bu, əsaslandırıla bilər. Yuxarıdakı misalda 128 milyon elementdən ibarət cədvəldə istifadə əmsalını 0,25-də saxlayırıqsa, onda biz ona 32 milyondan çox element yerləşdirə bilmərik və qalan 96 milyon slot itiriləcək - hər cüt üçün 8 bayt. , 768 MB itirilmiş yaddaş.
Nəzərə alın ki, söhbət sistem yaddaşından daha qiymətli resurs olan videokart yaddaşının itirilməsindən gedir. CUDA-nı dəstəkləyən müasir masaüstü qrafik kartlarının əksəriyyətinin ən azı 4 GB yaddaşı olsa da (yazı zamanı NVIDIA 2080 Ti-də 11 GB var), bu cür məbləğləri itirmək hələ də ən ağıllı qərar deyil.
Daha sonra mən zondlama dərinliyi ilə bağlı problemi olmayan video kartlar üçün hash cədvəllərinin yaradılması, eləcə də silinmiş yuvaların təkrar istifadəsi yolları haqqında daha çox yazacağam.
Səs dərinliyinin ölçülməsi
Açarın yoxlama dərinliyini müəyyən etmək üçün biz açarın hashini (onun ideal cədvəl indeksini) faktiki cədvəl indeksindən çıxara bilərik:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
İkinin ikisini tamamlayan ikili ədədlərin sehri və hash cədvəlinin tutumunun ikinin ikiyə bərabər olması səbəbindən bu yanaşma hətta açar indeks cədvəlin əvvəlinə köçürüldükdə də işləyəcək. Gəlin 1-ə həşlənmiş, lakin 3-cü yuvaya daxil edilmiş açarı götürək. Sonra 4 tutumu olan masa üçün biz alırıq. (3 — 1) & 3, bu 2-yə bərabərdir.
Nəticə
Suallarınız və ya şərhləriniz varsa, mənə eməktub göndərin və ya yeni mövzu açın.
Bu kod əla məqalələrdən ilham alaraq yazılmışdır:
Gələcəkdə video kartlar üçün hash masa tətbiqləri haqqında yazmağa və onların performansını təhlil etməyə davam edəcəyəm. Planlarıma GPU dostu olan məlumat strukturlarında atom əməliyyatlarından istifadə edərək zəncirləmə, Robin Qud hashing və kuku hashing daxildir.
Mənbə: www.habr.com