Ho vaig publicar a Github .
És una simple taula hash de GPU capaç de processar centenars de milions d'insercions per segon. Al meu portàtil NVIDIA GTX 1060, el codi insereix 64 milions de parells clau-valor generats aleatòriament en uns 210 ms i elimina 32 milions de parells en uns 64 ms.
És a dir, la velocitat d'un ordinador portàtil és d'aproximadament 300 milions d'insercions/s i 500 milions d'eliminacions/s.
La taula està escrita en CUDA, encara que la mateixa tècnica es pot aplicar a HLSL o GLSL. La implementació té diverses limitacions per garantir un alt rendiment en una targeta de vídeo:
- Només es processen claus de 32 bits i els mateixos valors.
- La taula hash té una mida fixa.
- I aquesta mida ha de ser igual a dos a la potència.
Per a claus i valors, cal reservar un marcador delimitador simple (al codi anterior és 0xffffffff).
Taula hash sense panys
La taula hash utilitza adreçament obert amb , és a dir, és simplement una matriu de parells clau-valor que s'emmagatzema a la memòria i té un rendiment de memòria cau superior. No es pot dir el mateix de l'encadenament, que consisteix a buscar un punter en una llista enllaçada. Una taula hash és una matriu simple que emmagatzema elements KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
La mida de la taula és una potència de dos, no un nombre primer, perquè una instrucció ràpida és suficient per aplicar la màscara pow2/AND, però l'operador de mòdul és molt més lent. Això és important en el cas del sondeig lineal, ja que en una cerca de taula lineal l'índex de ranura s'ha d'embolicar a cada ranura. I com a resultat, el cost de l'operació s'afegeix mòdul a cada ranura.
La taula només emmagatzema la clau i el valor de cada element, no un hash de la clau. Com que la taula només emmagatzema claus de 32 bits, el hash es calcula molt ràpidament. El codi anterior utilitza el hash Murmur3, que només realitza uns quants torns, XOR i multiplicacions.
La taula hash utilitza tècniques de protecció de bloqueig que són independents de l'ordre de la memòria. Fins i tot si algunes operacions d'escriptura pertorben l'ordre d'altres operacions d'aquest tipus, la taula hash encara mantindrà l'estat correcte. En parlarem a continuació. La tècnica funciona molt bé amb targetes de vídeo que executen milers de fils simultàniament.
Les claus i els valors de la taula hash s'inicien per buidar.
El codi es pot modificar per gestionar claus i valors de 64 bits també. Les claus requereixen operacions atòmiques de lectura, escriptura i comparació i intercanvi. I els valors requereixen operacions de lectura i escriptura atòmiques. Afortunadament, a CUDA, les operacions de lectura i escriptura per a valors de 32 i 64 bits són atòmiques sempre que estiguin alineades de manera natural (vegeu més avall). ), i les targetes de vídeo modernes admeten operacions de comparació i intercanvi atòmics de 64 bits. Per descomptat, en passar a 64 bits, el rendiment disminuirà lleugerament.
Estat de la taula hash
Cada parell clau-valor d'una taula hash pot tenir un d'aquests quatre estats:
- La clau i el valor estan buits. En aquest estat, la taula hash s'inicialitza.
- La clau s'ha anotat, però el valor encara no s'ha anotat. Si un altre fil està llegint dades, llavors torna buit. Això és normal, hauria passat el mateix si un altre fil d'execució hagués funcionat una mica abans, i estem parlant d'una estructura de dades concurrents.
- Es registren tant la clau com el valor.
- El valor està disponible per a altres fils d'execució, però la clau encara no ho està. Això pot passar perquè el model de programació CUDA té un model de memòria poc ordenat. Això és normal; en qualsevol cas, la clau encara està buida, encara que el valor ja no ho sigui.
Un matís important és que un cop s'ha escrit la clau a la ranura, ja no es mou, fins i tot si s'elimina la clau, en parlarem a continuació.
El codi de la taula hash fins i tot funciona amb models de memòria poc ordenats en els quals es desconeix l'ordre en què es llegeix i s'escriu la memòria. Quan mirem la inserció, la cerca i la supressió en una taula hash, recordeu que cada parell clau-valor es troba en un dels quatre estats descrits anteriorment.
Inserció en una taula hash
La funció CUDA que insereix parells clau-valor en una taula hash té aquest aspecte:
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);
}
}
Per inserir una clau, el codi itera a través de la matriu de la taula hash començant pel hash de la clau inserida. Cada ranura de la matriu realitza una operació de comparació i intercanvi atòmic que compara la clau d'aquesta ranura amb buida. Si es detecta una discrepància, la clau de la ranura s'actualitza amb la clau inserida i després es retorna la clau de la ranura original. Si aquesta clau original estava buida o coincidia amb la clau inserida, el codi va trobar una ranura adequada per inserir-la i va inserir el valor inserit a la ranura.
Si en una trucada del nucli gpu_hashtable_insert() hi ha diversos elements amb la mateixa clau, llavors qualsevol dels seus valors es pot escriure a la ranura de la clau. Això es considera normal: una de les escriptures clau-valor durant la trucada tindrà èxit, però com que tot això passa en paral·lel dins de diversos fils d'execució, no podem predir quina escriptura de memòria serà l'última.
Cerca de taula hash
Codi per cercar claus:
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);
}
}
Per trobar el valor d'una clau emmagatzemada en una taula, iterem per la matriu començant pel hash de la clau que estem buscant. A cada ranura, comprovem si la clau és la que busquem i, si és així, retornem el seu valor. També comprovem si la clau està buida, i si és així, avortem la cerca.
Si no podem trobar la clau, el codi retorna un valor buit.
Totes aquestes operacions de cerca es poden realitzar simultàniament mitjançant insercions i supressions. Cada parell de la taula tindrà un dels quatre estats descrits anteriorment per al flux.
Eliminació en una taula hash
Codi per esborrar claus:
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);
}
}
L'eliminació d'una clau es fa d'una manera inusual: deixem la clau a la taula i marquem el seu valor (no la clau en si) com a buida. Aquest codi és molt semblant a lookup(), excepte que quan es troba una coincidència en una clau, fa que el seu valor estigui buit.
Com s'ha esmentat anteriorment, una vegada que s'escriu una clau a una ranura, ja no es mou. Fins i tot quan s'elimina un element de la taula, la clau roman al seu lloc, el seu valor simplement queda buit. Això vol dir que no cal que utilitzem una operació d'escriptura atòmica per al valor de la ranura, perquè no importa si el valor actual està buit o no; encara quedarà buit.
Canviar la mida d'una taula hash
Podeu canviar la mida d'una taula hash creant una taula més gran i inserint-hi elements no buits de la taula antiga. No vaig implementar aquesta funcionalitat perquè volia mantenir el codi de mostra senzill. A més, als programes CUDA, l'assignació de memòria es fa sovint al codi amfitrió en lloc del nucli CUDA.
L’article descriu com modificar una estructura de dades protegida per bloqueig.
Competitivitat
Als fragments de codi de funció anteriors gpu_hashtable_insert(), _lookup() и _delete() processar un parell clau-valor alhora. I més baix gpu_hashtable_insert(), _lookup() и _delete() processar una matriu de parells en paral·lel, cada parell en un fil d'execució GPU independent:
// 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);
}
}
La taula hash resistent al bloqueig admet insercions, cerques i supressions concurrents. Com que els parells clau-valor sempre es troben en un dels quatre estats i les claus no es mouen, la taula garanteix la correcció fins i tot quan s'utilitzen diferents tipus d'operacions simultàniament.
Tanmateix, si processem un lot d'insercions i supressions en paral·lel, i si la matriu d'entrada de parells conté claus duplicades, aleshores no podrem predir quines parelles "guanyaran"; s'escriuran a la taula hash per últim. Suposem que anomenem el codi d'inserció amb una matriu d'entrada de parells A/0 B/1 A/2 C/3 A/4. Quan el codi s'ha completat, parelles B/1 и C/3 es garanteix la presència a la taula, però al mateix temps hi apareixerà qualsevol de les parelles A/0, A/2 o A/4. Això pot ser un problema o no; tot depèn de l'aplicació. És possible que sàpigues per endavant que no hi ha claus duplicades a la matriu d'entrada, o potser no t'importa quin valor es va escriure l'últim.
Si això és un problema per a vosaltres, haureu de separar els parells duplicats en diferents trucades al sistema CUDA. A CUDA, qualsevol operació que crida al nucli sempre es completa abans de la següent crida al nucli (almenys dins d'un fil. En diferents fils, els nuclis s'executen en paral·lel). A l'exemple anterior, si truqueu a un nucli amb A/0 B/1 A/2 C/3, i l'altre amb A/4, després la clau A obtindrà el valor 4.
Ara parlem de si les funcions haurien de ser lookup() и delete() utilitzeu un punter senzill o volàtil a una matriu de parells a la taula hash. Estats que:
El compilador pot optar per optimitzar les lectures i escriptures a la memòria global o compartida... Aquestes optimitzacions es poden desactivar mitjançant la paraula clau
volatile: ... qualsevol referència a aquesta variable es compila en una instrucció de lectura o escriptura de memòria real.
Les consideracions de correcció no requereixen aplicació volatile. Si el fil d'execució utilitza un valor en memòria cau d'una operació de lectura anterior, utilitzarà informació una mica obsoleta. Tot i així, aquesta és informació de l'estat correcte de la taula hash en un moment determinat de la trucada del nucli. Si necessiteu utilitzar la informació més recent, podeu utilitzar l'índex volatile, però aleshores el rendiment disminuirà lleugerament: segons les meves proves, en eliminar 32 milions d'elements, la velocitat va disminuir de 500 milions d'eliminacions/seg a 450 milions d'eliminacions/seg.
Productivitat
En la prova per inserir 64 milions d'elements i eliminar-ne 32 milions, la competència entre std::unordered_map i pràcticament no hi ha cap taula hash per a la GPU:
std::unordered_map va passar 70 ms inserint i eliminant elements i després alliberant-los unordered_map (Desfer-se de milions d'elements requereix molt de temps, perquè dins unordered_map es fan múltiples assignacions de memòria). Sincerament parlant, std:unordered_map restriccions completament diferents. És un únic fil d'execució de la CPU, admet valors-clau de qualsevol mida, funciona bé amb altes taxes d'utilització i mostra un rendiment estable després de múltiples supressions.
La durada de la taula hash per a la GPU i la comunicació entre programes va ser de 984 ms. Això inclou el temps dedicat a col·locar la taula a la memòria i suprimir-la (assignar 1 GB de memòria una vegada, que triga un temps a CUDA), inserir i suprimir elements i iterar sobre ells. També es tenen en compte totes les còpies cap a i des de la memòria de la targeta de vídeo.
La pròpia taula hash va trigar 271 ms a completar-se. Això inclou el temps que passa la targeta de vídeo inserint i suprimint elements, i no té en compte el temps dedicat a copiar a la memòria i a iterar sobre la taula resultant. Si la taula GPU viu durant molt de temps, o si la taula hash està continguda completament a la memòria de la targeta de vídeo (per exemple, per crear una taula hash que serà utilitzada per un altre codi GPU i no pel processador central), aleshores el resultat de la prova és rellevant.
La taula hash per a una targeta de vídeo demostra un alt rendiment a causa de l'alt rendiment i la paral·lelització activa.
Limitacions
L'arquitectura de la taula hash té alguns problemes que cal tenir en compte:
- El sondeig lineal es veu obstaculitzat per l'agrupació, la qual cosa fa que les claus de la taula no es col·loquin perfectament.
- Les claus no s'eliminen amb la funció
deletei amb el pas del temps desordena la taula.
Com a resultat, el rendiment d'una taula hash es pot degradar gradualment, sobretot si existeix durant molt de temps i té nombroses insercions i supressions. Una manera de mitigar aquests inconvenients és tornar a fer una taula nova amb una taxa d'utilització força baixa i filtrar les claus eliminades durant la repetició.
Per il·lustrar els problemes descrits, utilitzaré el codi anterior per crear una taula amb 128 milions d'elements i fer un bucle a través de 4 milions d'elements fins que hagi omplert 124 milions d'espais (taxa d'utilització d'uns 0,96). Aquí teniu la taula de resultats, cada fila és una crida al nucli CUDA per inserir 4 milions d'elements nous en una taula hash:
Taxa d'ús
Durada d'inserció 4 elements
0,00
11,608448 ms (361,314798 milions de claus/s)
0,03
11,751424 ms (356,918799 milions de claus/s)
0,06
11,942592 ms (351,205515 milions de claus/s)
0,09
12,081120 ms (347,178429 milions de claus/s)
0,12
12,242560 ms (342,600233 milions de claus/s)
0,16
12,396448 ms (338,347235 milions de claus/s)
0,19
12,533024 ms (334,660176 milions de claus/s)
0,22
12,703328 ms (330,173626 milions de claus/s)
0,25
12,884512 ms (325,530693 milions de claus/s)
0,28
13,033472 ms (321,810182 milions de claus/s)
0,31
13,239296 ms (316,807174 milions de claus/s)
0,34
13,392448 ms (313,184256 milions de claus/s)
0,37
13,624000 ms (307,861434 milions de claus/s)
0,41
13,875520 ms (302,280855 milions de claus/s)
0,44
14,126528 ms (296,909756 milions de claus/s)
0,47
14,399328 ms (291,284699 milions de claus/s)
0,50
14,690304 ms (285,515123 milions de claus/s)
0,53
15,039136 ms (278,892623 milions de claus/s)
0,56
15,478656 ms (270,973402 milions de claus/s)
0,59
15,985664 ms (262,379092 milions de claus/s)
0,62
16,668673 ms (251,627968 milions de claus/s)
0,66
17,587200 ms (238,486174 milions de claus/s)
0,69
18,690048 ms (224,413765 milions de claus/s)
0,72
20,278816 ms (206,831789 milions de claus/s)
0,75
22,545408 ms (186,038058 milions de claus/s)
0,78
26,053312 ms (160,989275 milions de claus/s)
0,81
31,895008 ms (131,503463 milions de claus/s)
0,84
42,103294 ms (99,619378 milions de claus/s)
0,87
61,849056 ms (67,815164 milions de claus/s)
0,90
105,695999 ms (39,682713 milions de claus/s)
0,94
240,204636 ms (17,461378 milions de claus/s)
A mesura que augmenta la utilització, el rendiment disminueix. Això no és desitjable en la majoria dels casos. Si una aplicació insereix elements en una taula i després els descarta (per exemple, quan es compta amb paraules en un llibre), això no és un problema. Però si l'aplicació utilitza una taula hash de llarga vida (per exemple, en un editor de gràfics per emmagatzemar parts no buides d'imatges on l'usuari insereix i suprimeix informació amb freqüència), aquest comportament pot ser problemàtic.
I va mesurar la profunditat de sondeig de la taula hash després de 64 milions d'insercions (factor d'utilització 0,5). La profunditat mitjana era de 0,4774, de manera que la majoria de claus es trobaven a la millor ranura possible o a una ranura de distància de la millor posició. La profunditat màxima de sondeig va ser de 60.
A continuació, vaig mesurar la profunditat de sondeig en una taula amb 124 milions d'insercions (factor d'utilització 0,97). La profunditat mitjana ja era de 10,1757, i la màxima - 6474 (!!). El rendiment de la detecció lineal disminueix significativament a taxes d'utilització elevades.
El millor és mantenir baixa la taxa d'utilització d'aquesta taula hash. Però després augmentem el rendiment a costa del consum de memòria. Afortunadament, en el cas de claus i valors de 32 bits, això es pot justificar. Si a l'exemple anterior, en una taula amb 128 milions d'elements, mantenim el factor d'utilització de 0,25, llavors no hi podem col·locar més de 32 milions d'elements i es perdran els 96 milions de ranures restants: 8 bytes per cada parell. , 768 MB de memòria perduda.
Tingueu en compte que estem parlant de la pèrdua de memòria de la targeta de vídeo, que és un recurs més valuós que la memòria del sistema. Tot i que la majoria de les targetes gràfiques d'escriptori modernes que admeten CUDA tenen almenys 4 GB de memòria (en el moment d'escriure, la NVIDIA 2080 Ti en té 11 GB), encara no seria la decisió més sàvia perdre aquestes quantitats.
Més endavant escriuré més sobre la creació de taules hash per a targetes de vídeo que no tenen problemes amb la profunditat de sondeig, així com maneres de reutilitzar les ranures suprimides.
Mesura de la profunditat del sondeig
Per determinar la profunditat de sondeig d'una clau, podem extreure el hash de la clau (el seu índex de taula ideal) del seu índex de taula real:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
A causa de la màgia dels números binaris del complement a dos de dos i del fet que la capacitat de la taula hash és de dos a la potència de dos, aquest enfocament funcionarà fins i tot quan l'índex clau es mou al principi de la taula. Agafem una clau amb hash a 1, però que s'insereix a la ranura 3. Aleshores, per a una taula amb capacitat 4, obtenim (3 — 1) & 3, que equival a 2.
Conclusió
Si teniu preguntes o comentaris, envieu-me un correu electrònic a o obre un tema nou a.
Aquest codi es va escriure inspirant-se en articles excel·lents:
En el futur, continuaré escrivint sobre implementacions de taules hash per a targetes de vídeo i analitzaré el seu rendiment. Els meus plans inclouen encadenar, hash Robin Hood i hash cucut mitjançant operacions atòmiques en estructures de dades compatibles amb la GPU.
Font: www.habr.com