L-am postat pe Github .
Este un simplu tabel hash GPU capabil să proceseze sute de milioane de inserări pe secundă. Pe laptopul meu NVIDIA GTX 1060, codul inserează 64 de milioane de perechi cheie-valoare generate aleatoriu în aproximativ 210 ms și elimină 32 de milioane de perechi în aproximativ 64 ms.
Adică, viteza pe un laptop este de aproximativ 300 de milioane de inserări/sec și 500 de milioane de ștergeri/sec.
Tabelul este scris în CUDA, deși aceeași tehnică poate fi aplicată la HLSL sau GLSL. Implementarea are câteva limitări pentru a asigura performanțe ridicate pe o placă video:
- Sunt procesate doar cheile pe 32 de biți și aceleași valori.
- Tabelul hash are o dimensiune fixă.
- Și această dimensiune trebuie să fie egală cu două la putere.
Pentru chei și valori, trebuie să rezervați un marker delimitator simplu (în codul de mai sus acesta este 0xffffffff).
Tabel de hash fără încuietori
Tabelul hash folosește adresarea deschisă cu , adică este pur și simplu o matrice de perechi cheie-valoare care este stocată în memorie și are performanțe superioare în cache. Nu același lucru se poate spune despre înlănțuire, care implică căutarea unui pointer într-o listă legată. Un tabel hash este o matrice simplă care stochează elemente KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Mărimea tabelului este o putere de doi, nu un număr prim, deoarece o instrucțiune rapidă este suficientă pentru a aplica masca pow2/AND, dar operatorul de modul este mult mai lent. Acest lucru este important în cazul sondării liniare, deoarece într-o căutare a unui tabel liniar indexul slotului trebuie să fie înfășurat în fiecare slot. Și, ca rezultat, costul operațiunii este adăugat modulo în fiecare slot.
Tabelul stochează doar cheia și valoarea pentru fiecare element, nu un hash al cheii. Deoarece tabelul stochează doar chei pe 32 de biți, hash-ul este calculat foarte rapid. Codul de mai sus folosește hash-ul Murmur3, care efectuează doar câteva ture, XOR-uri și înmulțiri.
Tabelul hash folosește tehnici de protecție de blocare care sunt independente de ordinea memoriei. Chiar dacă unele operații de scriere perturbă ordinea altor astfel de operațiuni, tabelul hash va menține în continuare starea corectă. Despre asta vom vorbi mai jos. Tehnica funcționează excelent cu plăcile video care rulează mii de fire simultan.
Cheile și valorile din tabelul hash sunt inițializate pentru a se goli.
Codul poate fi modificat pentru a gestiona, de asemenea, cheile și valorile pe 64 de biți. Cheile necesită operații atomice de citire, scriere și comparare și schimbare. Și valorile necesită operații atomice de citire și scriere. Din fericire, în CUDA, operațiunile de citire-scriere pentru valorile pe 32 și 64 de biți sunt atomice atâta timp cât sunt aliniate în mod natural (vezi mai jos). ), iar plăcile video moderne acceptă operațiuni de comparare și schimb atomic pe 64 de biți. Desigur, când treceți la 64 de biți, performanța va scădea ușor.
Starea tabelului hash
Fiecare pereche cheie-valoare dintr-un tabel hash poate avea una dintre cele patru stări:
- Cheia și valoarea sunt goale. În această stare, tabelul hash este inițializat.
- Cheia a fost scrisă, dar valoarea nu a fost încă scrisă. Dacă un alt fir de execuție citește în prezent date, acesta revine gol. Este normal, același lucru s-ar fi întâmplat dacă un alt thread de execuție ar fi funcționat puțin mai devreme și vorbim despre o structură de date concurentă.
- Atât cheia, cât și valoarea sunt înregistrate.
- Valoarea este disponibilă pentru alte fire de execuție, dar cheia nu este încă. Acest lucru se poate întâmpla deoarece modelul de programare CUDA are un model de memorie ordonat vag. Acest lucru este normal; în orice caz, cheia este încă goală, chiar dacă valoarea nu mai este așa.
O nuanță importantă este că odată ce cheia a fost scrisă în slot, aceasta nu se mai mișcă - chiar dacă cheia este ștearsă, despre asta vom vorbi mai jos.
Codul tabelului hash funcționează chiar și cu modele de memorie ordonate vag, în care ordinea în care este citită și scrisă memoria este necunoscută. Pe măsură ce ne uităm la inserarea, căutarea și ștergerea într-un tabel hash, rețineți că fiecare pereche cheie-valoare se află în una dintre cele patru stări descrise mai sus.
Inserarea într-un tabel hash
Funcția CUDA care inserează perechi cheie-valoare într-un tabel hash arată astfel:
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);
}
}
Pentru a insera o cheie, codul iterează prin matricea tabelului hash, începând cu hash-ul cheii introduse. Fiecare slot din matrice efectuează o operație de comparare și schimbare atomică care compară cheia din acel slot cu golirea. Dacă este detectată o nepotrivire, cheia din slot este actualizată cu cheia introdusă și apoi cheia originală a slotului este returnată. Dacă această cheie originală era goală sau se potrivea cu cheia introdusă, atunci codul a găsit un slot potrivit pentru inserare și a introdus valoarea introdusă în slot.
Dacă într-un apel de nucleu gpu_hashtable_insert() există mai multe elemente cu aceeași cheie, apoi oricare dintre valorile lor poate fi scrisă în slotul cheii. Acest lucru este considerat normal: una dintre scrierile cheie-valoare în timpul apelului va reuși, dar deoarece toate acestea se întâmplă în paralel în mai multe fire de execuție, nu putem prezice care scriere în memorie va fi ultima.
Căutare tabel hash
Cod pentru căutarea cheilor:
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);
}
}
Pentru a găsi valoarea unei chei stocate într-un tabel, iterăm prin matrice începând cu hash-ul cheii pe care o căutăm. În fiecare slot, verificăm dacă cheia este cea pe care o căutăm, iar dacă da, îi returnăm valoarea. De asemenea, verificăm dacă cheia este goală, iar dacă da, anulăm căutarea.
Dacă nu putem găsi cheia, codul returnează o valoare goală.
Toate aceste operațiuni de căutare pot fi efectuate concomitent prin inserări și ștergeri. Fiecare pereche din tabel va avea una dintre cele patru stări descrise mai sus pentru flux.
Ștergerea într-un tabel hash
Cod pentru ștergerea cheilor:
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);
}
}
Ștergerea unei chei se face într-un mod neobișnuit: lăsăm cheia în tabel și marcăm valoarea acesteia (nu cheia în sine) ca goală. Acest cod este foarte asemănător cu lookup(), cu excepția faptului că atunci când se găsește o potrivire pe o cheie, aceasta își face valoarea goală.
După cum am menționat mai sus, odată ce o cheie este scrisă într-un slot, aceasta nu mai este mutată. Chiar și atunci când un element este șters din tabel, cheia rămâne pe loc, valoarea sa devine pur și simplu goală. Aceasta înseamnă că nu trebuie să folosim o operație de scriere atomică pentru valoarea slotului, deoarece nu contează dacă valoarea curentă este goală sau nu - va deveni în continuare goală.
Redimensionarea unui tabel hash
Puteți modifica dimensiunea unui tabel hash creând un tabel mai mare și inserând elemente negoale din tabelul vechi în el. Nu am implementat această funcționalitate pentru că am vrut să păstrez exemplul de cod simplu. Mai mult, în programele CUDA, alocarea memoriei se face adesea în codul gazdă mai degrabă decât în nucleul CUDA.
Articolul descrie cum se modifică o astfel de structură de date protejată de blocare.
Competitivitate
În fragmentele de cod al funcției de mai sus gpu_hashtable_insert(), _lookup() и _delete() procesează câte o pereche cheie-valoare la un moment dat. Și mai jos gpu_hashtable_insert(), _lookup() и _delete() procesează o serie de perechi în paralel, fiecare pereche într-un fir de execuție GPU separat:
// 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);
}
}
Tabelul hash rezistent la blocare acceptă inserări, căutări și ștergeri simultane. Deoarece perechile cheie-valoare sunt întotdeauna în una din cele patru stări și cheile nu se mișcă, tabelul garantează corectitudinea chiar și atunci când diferite tipuri de operații sunt utilizate simultan.
Cu toate acestea, dacă procesăm un lot de inserări și ștergeri în paralel și dacă matricea de perechi de intrare conține chei duplicate, atunci nu vom putea prezice ce perechi vor „câștiga” – vor fi scrise ultima în tabelul hash. Să presupunem că am numit codul de inserare cu un tablou de intrare de perechi A/0 B/1 A/2 C/3 A/4. Când codul se completează, perechi B/1 и C/3 sunt garantate a fi prezente în tabel, dar, în același timp, oricare dintre perechi va apărea în acesta A/0, A/2 sau A/4. Aceasta poate fi sau nu o problemă - totul depinde de aplicație. Este posibil să știți dinainte că nu există chei duplicate în matricea de intrare sau s-ar putea să nu vă pese care valoare a fost scrisă ultima.
Dacă aceasta este o problemă pentru dvs., atunci trebuie să separați perechile duplicate în diferite apeluri de sistem CUDA. În CUDA, orice operație care apelează nucleul se finalizează întotdeauna înainte de următorul apel kernel (cel puțin într-un fir. În diferite fire de execuție, nucleele sunt executate în paralel). În exemplul de mai sus, dacă apelați un nucleu cu A/0 B/1 A/2 C/3, iar celălalt cu A/4, apoi cheia A va primi valoarea 4.
Acum să vorbim despre dacă funcțiile ar trebui lookup() и delete() utilizați un indicator simplu sau volatil către o matrice de perechi din tabelul hash. Afirmă că:
Compilatorul poate alege să optimizeze citirile și scrierile în memoria globală sau partajată... Aceste optimizări pot fi dezactivate folosind cuvântul cheie
volatile: ... orice referință la această variabilă este compilată într-o instrucțiune de citire sau scriere în memorie reală.
Considerațiile de corectitudine nu necesită aplicare volatile. Dacă firul de execuție folosește o valoare stocată în cache dintr-o operație de citire anterioară, atunci va folosi informații ușor învechite. Dar totuși, acestea sunt informații din starea corectă a tabelului hash la un anumit moment al apelului nucleului. Dacă trebuie să utilizați cele mai recente informații, puteți utiliza indexul volatile, dar apoi performanța va scădea ușor: conform testelor mele, la ștergerea a 32 de milioane de elemente, viteza a scăzut de la 500 de milioane de ștergeri/sec la 450 de milioane de ștergeri/sec.
productivitate
În testul pentru introducerea a 64 de milioane de elemente și ștergerea a 32 de milioane dintre ele, concurența între std::unordered_map și practic nu există nicio tabelă hash pentru GPU:
std::unordered_map a petrecut 70 ms inserând și scoțând elemente și apoi eliberându-le unordered_map (a scăpa de milioane de elemente necesită mult timp, pentru că în interior unordered_map se fac alocări multiple de memorie). Sincer vorbind, std:unordered_map restricții complet diferite. Este un singur fir de execuție al procesorului, acceptă valori cheie de orice dimensiune, funcționează bine la rate mari de utilizare și arată performanțe stabile după ștergeri multiple.
Durata tabelului hash pentru GPU și comunicarea între programe a fost de 984 ms. Aceasta include timpul petrecut pentru plasarea tabelului în memorie și ștergerea acestuia (alocarea a 1 GB de memorie o dată, ceea ce durează ceva timp în CUDA), inserarea și ștergerea elementelor și iterarea peste ele. De asemenea, sunt luate în considerare toate copiile către și din memoria plăcii video.
Tabelul hash în sine a durat 271 ms. Aceasta include timpul petrecut de placa video inserând și ștergând elemente și nu ia în considerare timpul petrecut cu copierea în memorie și iterarea peste tabelul rezultat. Dacă tabelul GPU trăiește mult timp sau dacă tabelul hash este conținut în întregime în memoria plăcii video (de exemplu, pentru a crea un tabel hash care va fi folosit de alt cod GPU și nu de procesorul central), atunci rezultatul testului este relevant.
Tabelul hash pentru o placă video demonstrează performanțe ridicate datorită debitului mare și paralelizării active.
Limitări
Arhitectura tabelului hash are câteva probleme de care trebuie să fii conștient:
- Sondarea liniară este împiedicată de grupare, ceea ce face ca cheile din tabel să fie plasate mai puțin decât perfect.
- Cheile nu sunt eliminate folosind funcția
deleteiar de-a lungul timpului aglomerează masa.
Drept urmare, performanța unui tabel hash se poate degrada treptat, mai ales dacă există de mult timp și are numeroase inserări și ștergeri. O modalitate de a atenua aceste dezavantaje este să rehașați într-un tabel nou cu o rată de utilizare destul de scăzută și să filtrați cheile eliminate în timpul rehașării.
Pentru a ilustra problemele descrise, voi folosi codul de mai sus pentru a crea un tabel cu 128 de milioane de elemente și voi trece în buclă prin 4 milioane de elemente până când voi umple 124 de milioane de sloturi (rată de utilizare de aproximativ 0,96). Iată tabelul cu rezultate, fiecare rând este un apel de nucleu CUDA pentru a insera 4 milioane de elemente noi într-un singur tabel hash:
Rata de folosire
Durata de inserare 4 elemente
0,00
11,608448 ms (361,314798 milioane de taste/sec.)
0,03
11,751424 ms (356,918799 milioane de taste/sec.)
0,06
11,942592 ms (351,205515 milioane de taste/sec.)
0,09
12,081120 ms (347,178429 milioane de taste/sec.)
0,12
12,242560 ms (342,600233 milioane de taste/sec.)
0,16
12,396448 ms (338,347235 milioane de taste/sec.)
0,19
12,533024 ms (334,660176 milioane de taste/sec.)
0,22
12,703328 ms (330,173626 milioane de taste/sec.)
0,25
12,884512 ms (325,530693 milioane de taste/sec.)
0,28
13,033472 ms (321,810182 milioane de taste/sec.)
0,31
13,239296 ms (316,807174 milioane de taste/sec.)
0,34
13,392448 ms (313,184256 milioane de taste/sec.)
0,37
13,624000 ms (307,861434 milioane de taste/sec.)
0,41
13,875520 ms (302,280855 milioane de taste/sec.)
0,44
14,126528 ms (296,909756 milioane de taste/sec.)
0,47
14,399328 ms (291,284699 milioane de taste/sec.)
0,50
14,690304 ms (285,515123 milioane de taste/sec.)
0,53
15,039136 ms (278,892623 milioane de taste/sec.)
0,56
15,478656 ms (270,973402 milioane de taste/sec.)
0,59
15,985664 ms (262,379092 milioane de taste/sec.)
0,62
16,668673 ms (251,627968 milioane de taste/sec.)
0,66
17,587200 ms (238,486174 milioane de taste/sec.)
0,69
18,690048 ms (224,413765 milioane de taste/sec.)
0,72
20,278816 ms (206,831789 milioane de taste/sec.)
0,75
22,545408 ms (186,038058 milioane de taste/sec.)
0,78
26,053312 ms (160,989275 milioane de taste/sec.)
0,81
31,895008 ms (131,503463 milioane de taste/sec.)
0,84
42,103294 ms (99,619378 milioane de taste/sec.)
0,87
61,849056 ms (67,815164 milioane de taste/sec.)
0,90
105,695999 ms (39,682713 milioane de taste/sec.)
0,94
240,204636 ms (17,461378 milioane de taste/sec.)
Pe măsură ce utilizarea crește, performanța scade. Acest lucru nu este de dorit în majoritatea cazurilor. Dacă o aplicație inserează elemente într-un tabel și apoi le aruncă (de exemplu, când numără cuvintele dintr-o carte), atunci aceasta nu este o problemă. Dar dacă aplicația folosește un tabel hash de lungă durată (de exemplu, într-un editor grafic pentru a stoca părți nevide ale imaginilor în care utilizatorul inserează și șterge frecvent informații), atunci acest comportament poate fi problematic.
Și am măsurat adâncimea de sondare a tabelului hash după 64 de milioane de inserții (factor de utilizare 0,5). Adâncimea medie a fost de 0,4774, așa că majoritatea cheilor se aflau fie în cel mai bun slot posibil, fie la un slot distanță de cea mai bună poziție. Adâncimea maximă de sondare a fost de 60.
Apoi am măsurat adâncimea de sondare pe o masă cu 124 de milioane de inserții (factor de utilizare 0,97). Adâncimea medie era deja de 10,1757, iar maxima - 6474 (!!). Performanța de detectare liniară scade semnificativ la rate mari de utilizare.
Cel mai bine este să mențineți rata de utilizare a acestui tabel hash scăzut. Dar apoi creștem performanța în detrimentul consumului de memorie. Din fericire, în cazul cheilor și valorilor pe 32 de biți, acest lucru poate fi justificat. Dacă în exemplul de mai sus, într-un tabel cu 128 de milioane de elemente, păstrăm factorul de utilizare de 0,25, atunci nu putem plasa mai mult de 32 de milioane de elemente în el, iar restul de 96 de milioane de sloturi se vor pierde - 8 octeți pentru fiecare pereche , 768 MB de memorie pierdută.
Vă rugăm să rețineți că vorbim despre pierderea memoriei plăcii video, care este o resursă mai valoroasă decât memoria de sistem. Deși majoritatea plăcilor grafice desktop moderne care acceptă CUDA au cel puțin 4 GB de memorie (la momentul scrierii, NVIDIA 2080 Ti are 11 GB), tot nu ar fi cea mai înțeleaptă decizie să pierzi astfel de sume.
Mai târziu voi scrie mai multe despre crearea tabelelor hash pentru plăcile video care nu au probleme cu adâncimea de sondare, precum și modalități de reutilizare a sloturilor șterse.
Măsurarea adâncimii de sunet
Pentru a determina adâncimea de sondare a unei chei, putem extrage hash-ul cheii (indexul său ideal de tabel) din indexul său real de tabel:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Datorită magiei numerelor binare în complement a doi și a faptului că capacitatea tabelului hash este de două la puterea lui doi, această abordare va funcționa chiar și atunci când indexul cheie este mutat la începutul tabelului. Să luăm o cheie care are hash la 1, dar este introdusă în slotul 3. Apoi, pentru un tabel cu capacitate 4, obținem (3 — 1) & 3, care este echivalent cu 2.
Concluzie
Dacă aveți întrebări sau comentarii, vă rugăm să-mi trimiteți un e-mail la adresa sau deschideți un subiect nou în.
Acest cod a fost scris inspirat de articole excelente:
În viitor, voi continua să scriu despre implementările tabelelor hash pentru plăcile video și voi analiza performanța acestora. Planurile mele includ înlănțuire, hashing Robin Hood și hashing cuc folosind operațiuni atomice în structuri de date care sunt prietenoase cu GPU.
Sursa: www.habr.com