Ik heb het op Github geplaatst .
Het is een eenvoudige GPU-hashtabel die honderden miljoenen inserts per seconde kan verwerken. Op mijn NVIDIA GTX 1060-laptop voegt de code 64 miljoen willekeurig gegenereerde sleutel-waardeparen in in ongeveer 210 ms en verwijdert 32 miljoen paren in ongeveer 64 ms.
Dat wil zeggen, de snelheid op een laptop is ongeveer 300 miljoen invoegingen/sec en 500 miljoen verwijderingen/sec.
De tabel is geschreven in CUDA, hoewel dezelfde techniek kan worden toegepast op HLSL of GLSL. De implementatie heeft verschillende beperkingen om hoge prestaties op een videokaart te garanderen:
- Er worden alleen 32-bits sleutels en dezelfde waarden verwerkt.
- De hashtabel heeft een vaste maat.
- En deze grootte moet gelijk zijn aan twee tot de macht.
Voor sleutels en waarden moet u een eenvoudig scheidingsteken reserveren (in de bovenstaande code is dit 0xffffffff).
Hasjtafel zonder sloten
De hashtabel maakt gebruik van open adressering met , dat wil zeggen dat het eenvoudigweg een reeks sleutel-waardeparen is die in het geheugen wordt opgeslagen en superieure cacheprestaties heeft. Hetzelfde kan niet gezegd worden voor chaining, waarbij gezocht wordt naar een pointer in een gekoppelde lijst. Een hashtabel is een eenvoudige array waarin elementen worden opgeslagen KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
De grootte van de tabel is een macht van twee, geen priemgetal, omdat één snelle instructie voldoende is om het pow2/AND-masker toe te passen, maar de modulus-operator is veel langzamer. Dit is belangrijk in het geval van lineair onderzoek, omdat bij het opzoeken van een lineaire tabel de slotindex in elk slot moet worden opgenomen. En als gevolg daarvan worden de kosten van de operatie in elk slot modulo opgeteld.
De tabel slaat alleen de sleutel en waarde voor elk element op, geen hash van de sleutel. Omdat de tabel slechts 32-bits sleutels opslaat, wordt de hash zeer snel berekend. De bovenstaande code gebruikt de Murmur3-hash, die slechts enkele verschuivingen, XOR's en vermenigvuldigingen uitvoert.
De hashtabel maakt gebruik van vergrendelingsbeschermingstechnieken die onafhankelijk zijn van de geheugenvolgorde. Zelfs als sommige schrijfbewerkingen de volgorde van andere dergelijke bewerkingen verstoren, behoudt de hashtabel nog steeds de juiste status. We zullen hier hieronder over praten. De techniek werkt uitstekend met videokaarten die duizenden threads tegelijk uitvoeren.
De sleutels en waarden in de hashtabel zijn geïnitialiseerd op leeg.
De code kan worden aangepast om ook 64-bits sleutels en waarden te verwerken. Sleutels vereisen atomaire lees-, schrijf- en vergelijkings- en wisselbewerkingen. En waarden vereisen atomaire lees- en schrijfbewerkingen. Gelukkig zijn in CUDA lees-schrijfbewerkingen voor 32- en 64-bits waarden atomair, zolang ze op natuurlijke wijze uitgelijnd zijn (zie hieronder). ), en moderne videokaarten ondersteunen 64-bit atomaire vergelijkings- en uitwisselingsbewerkingen. Bij de overstap naar 64 bits zullen de prestaties uiteraard iets afnemen.
Status van de hashtabel
Elk sleutel-waardepaar in een hashtabel kan een van de vier statussen hebben:
- Sleutel en waarde zijn leeg. In deze toestand wordt de hashtabel geïnitialiseerd.
- De sleutel is opgeschreven, maar de waarde is nog niet geschreven. Als een andere thread momenteel gegevens leest, keert deze leeg terug. Dit is normaal, hetzelfde zou zijn gebeurd als een andere uitvoeringsdraad iets eerder had gewerkt, en we hebben het over een gelijktijdige datastructuur.
- Zowel de sleutel als de waarde worden vastgelegd.
- De waarde is beschikbaar voor andere uitvoeringsthreads, maar de sleutel nog niet. Dit kan gebeuren omdat het CUDA-programmeermodel een losjes geordend geheugenmodel heeft. Dit is normaal; in ieder geval is de sleutel nog steeds leeg, ook al is de waarde niet langer leeg.
Een belangrijke nuance is dat zodra de sleutel in het slot is geschreven, deze niet meer beweegt - zelfs als de sleutel wordt verwijderd, zullen we hier hieronder over praten.
De hashtabelcode werkt zelfs met los geordende geheugenmodellen waarin de volgorde waarin geheugen wordt gelezen en geschreven onbekend is. Als we kijken naar het invoegen, opzoeken en verwijderen in een hashtabel, onthoud dan dat elk sleutel-waardepaar zich in een van de vier hierboven beschreven toestanden bevindt.
Invoegen in een hashtabel
De CUDA-functie die sleutel-waardeparen in een hashtabel invoegt, ziet er als volgt uit:
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);
}
}
Om een sleutel in te voegen, herhaalt de code de hashtabelarray, te beginnen met de hash van de ingevoegde sleutel. Elke sleuf in de array voert een atomaire vergelijk-en-wisseloperatie uit, waarbij de sleutel in die sleuf wordt vergeleken met leeg. Als er een verkeerde combinatie wordt gedetecteerd, wordt de sleutel in het slot bijgewerkt met de ingestoken sleutel en wordt vervolgens de originele slotsleutel teruggegeven. Als deze originele sleutel leeg was of overeenkwam met de ingestoken sleutel, vond de code een geschikt slot om in te steken en plaatste de ingevoegde waarde in het slot.
Als het in één kerneloproep is gpu_hashtable_insert() er zijn meerdere elementen met dezelfde sleutel, dan kan elk van hun waarden naar het sleutelslot worden geschreven. Dit wordt als normaal beschouwd: een van de sleutelwaardeschrijfbewerkingen tijdens de aanroep zal slagen, maar aangezien dit allemaal parallel gebeurt binnen verschillende uitvoeringsthreads, kunnen we niet voorspellen welke geheugenschrijfbewerking de laatste zal zijn.
Hashtabel opzoeken
Code voor het zoeken naar sleutels:
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);
}
}
Om de waarde te vinden van een sleutel die in een tabel is opgeslagen, doorlopen we de array, beginnend met de hash van de sleutel die we zoeken. In elk slot controleren we of de sleutel de sleutel is die we zoeken, en zo ja, dan geven we de waarde ervan terug. Ook controleren wij of de sleutel leeg is, en zo ja, dan breken wij de zoektocht af.
Als we de sleutel niet kunnen vinden, retourneert de code een lege waarde.
Al deze zoekbewerkingen kunnen gelijktijdig worden uitgevoerd via invoegingen en verwijderingen. Elk paar in de tabel heeft een van de vier hierboven beschreven toestanden voor de stroom.
Verwijderen in een hashtabel
Code voor het verwijderen van sleutels:
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);
}
}
Het verwijderen van een sleutel gebeurt op een ongebruikelijke manier: we laten de sleutel in de tabel staan en markeren de waarde ervan (niet de sleutel zelf) als leeg. Deze code lijkt erg op lookup(), behalve dat wanneer er een overeenkomst wordt gevonden op een sleutel, de waarde ervan leeg wordt gemaakt.
Zoals hierboven vermeld, wordt een sleutel die eenmaal naar een slot is geschreven, niet meer verplaatst. Zelfs als een element uit de tabel wordt verwijderd, blijft de sleutel op zijn plaats; de waarde ervan wordt eenvoudigweg leeg. Dit betekent dat we geen atomaire schrijfbewerking voor de slotwaarde hoeven te gebruiken, omdat het niet uitmaakt of de huidige waarde leeg is of niet - deze zal nog steeds leeg worden.
Het formaat van een hashtabel wijzigen
U kunt de grootte van een hashtabel wijzigen door een grotere tabel te maken en daarin niet-lege elementen uit de oude tabel in te voegen. Ik heb deze functionaliteit niet geïmplementeerd omdat ik de voorbeeldcode eenvoudig wilde houden. Bovendien wordt geheugentoewijzing in CUDA-programma's vaak gedaan in de hostcode in plaats van in de CUDA-kernel.
In het artikel beschrijft hoe u een dergelijke met een slot beveiligde datastructuur kunt wijzigen.
Concurrentievermogen
In de bovenstaande functiecodefragmenten gpu_hashtable_insert(), _lookup() и _delete() één sleutel-waardepaar tegelijk verwerken. En lager gpu_hashtable_insert(), _lookup() и _delete() een reeks paren parallel verwerken, elk paar in een afzonderlijke GPU-uitvoeringsthread:
// 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);
}
}
De slotbestendige hashtabel ondersteunt gelijktijdige invoegingen, opzoekingen en verwijderingen. Omdat sleutel-waardeparen zich altijd in een van de vier toestanden bevinden en de sleutels niet bewegen, garandeert de tabel juistheid, zelfs wanneer verschillende soorten bewerkingen tegelijkertijd worden gebruikt.
Als we echter een reeks invoegingen en verwijderingen parallel verwerken, en als de invoerarray van paren dubbele sleutels bevat, kunnen we niet voorspellen welke paren zullen “winnen” en als laatste naar de hashtabel worden geschreven. Laten we zeggen dat we de invoegcode hebben aangeroepen met een invoerarray van paren A/0 B/1 A/2 C/3 A/4. Wanneer de code is voltooid, paren B/1 и C/3 zijn gegarandeerd aanwezig in de tabel, maar tegelijkertijd zal een van de paren erin verschijnen A/0, A/2 of A/4. Dit kan wel of geen probleem zijn - het hangt allemaal af van de toepassing. Mogelijk weet u van tevoren dat er geen dubbele sleutels in de invoerarray voorkomen, of maakt het u misschien niet uit welke waarde het laatst is geschreven.
Als dit een probleem voor u is, moet u de dubbele paren scheiden in verschillende CUDA-systeemaanroepen. In CUDA wordt elke bewerking die de kernel aanroept altijd voltooid vóór de volgende kernelaanroep (tenminste binnen één thread. In verschillende threads worden kernels parallel uitgevoerd). In het bovenstaande voorbeeld, als je één kernel aanroept met A/0 B/1 A/2 C/3, en de andere met A/4en vervolgens de sleutel A de waarde zal krijgen 4.
Laten we het nu hebben over de vraag of functies dat wel zouden moeten doen lookup() и delete() gebruik een gewone of vluchtige verwijzing naar een reeks paren in de hashtabel. Zegt dat:
De compiler kan ervoor kiezen om het lezen en schrijven naar het globale of gedeelde geheugen te optimaliseren... Deze optimalisaties kunnen worden uitgeschakeld met behulp van het trefwoord
volatile: ... elke verwijzing naar deze variabele wordt gecompileerd in een echte geheugenlees- of schrijfinstructie.
Correctheidsoverwegingen behoeven geen toepassing volatile. Als de uitvoeringsthread een in de cache opgeslagen waarde uit een eerdere leesbewerking gebruikt, wordt er enigszins verouderde informatie gebruikt. Maar toch is dit informatie over de juiste status van de hashtabel op een bepaald moment van de kernelaanroep. Als u de meest recente informatie nodig heeft, kunt u de index gebruiken volatile, maar dan zullen de prestaties iets afnemen: volgens mijn tests daalde de snelheid bij het verwijderen van 32 miljoen elementen van 500 miljoen verwijderingen/sec naar 450 miljoen verwijderingen/sec.
Производительность
In de test voor het invoegen van 64 miljoen elementen en het verwijderen van 32 miljoen daarvan, concurrentie tussen std::unordered_map en er is vrijwel geen hashtabel voor de GPU:
std::unordered_map besteedde 70 ms aan het invoegen en verwijderen van elementen en het vervolgens vrijmaken ervan unordered_map (het verwijderen van miljoenen elementen kost veel tijd, omdat inside unordered_map er worden meerdere geheugentoewijzingen uitgevoerd). Eerlijk gezegd, std:unordered_map totaal verschillende beperkingen. Het is een enkele CPU-uitvoeringsthread, ondersteunt sleutelwaarden van elke grootte, presteert goed bij hoge bezettingsgraden en vertoont stabiele prestaties na meerdere verwijderingen.
De duur van de hashtabel voor de GPU en communicatie tussen programma's was 984 ms. Dit omvat de tijd die wordt besteed aan het plaatsen van de tabel in het geheugen en het verwijderen ervan (één keer 1 GB geheugen toewijzen, wat enige tijd kost in CUDA), het invoegen en verwijderen van elementen, en het herhalen ervan. Er wordt ook rekening gehouden met alle kopieën van en naar het videokaartgeheugen.
Het voltooien van de hashtabel zelf duurde 271 ms. Dit omvat de tijd die de videokaart besteedt aan het invoegen en verwijderen van elementen, en houdt geen rekening met de tijd die wordt besteed aan het kopiëren naar het geheugen en het herhalen van de resulterende tabel. Als de GPU-tabel lange tijd meegaat, of als de hashtabel volledig in het geheugen van de videokaart staat (bijvoorbeeld om een hashtabel te maken die door andere GPU-code en niet door de centrale processor wordt gebruikt), dan het testresultaat is relevant.
De hashtabel voor een videokaart vertoont hoge prestaties dankzij de hoge doorvoer en actieve parallellisatie.
Beperkingen
Er zijn een aantal problemen waarmee u rekening moet houden bij de hashtabelarchitectuur:
- Lineair sonderen wordt belemmerd door clustering, waardoor de toetsen in de tafel niet helemaal perfect worden geplaatst.
- Sleutels worden niet verwijderd met behulp van de functie
deleteen na verloop van tijd maken ze de tafel rommelig.
Als gevolg hiervan kunnen de prestaties van een hashtabel geleidelijk afnemen, vooral als deze al langere tijd bestaat en talloze invoegingen en verwijderingen bevat. Eén manier om deze nadelen te verzachten is door opnieuw te hashen naar een nieuwe tabel met een vrij lage bezettingsgraad en de verwijderde sleutels eruit te filteren tijdens het opnieuw hashen.
Om de beschreven problemen te illustreren, gebruik ik de bovenstaande code om een tabel met 128 miljoen elementen te maken en door 4 miljoen elementen te lopen totdat ik 124 miljoen slots heb gevuld (gebruikspercentage van ongeveer 0,96). Hier is de resultatentabel, elke rij is een CUDA-kerneloproep om 4 miljoen nieuwe elementen in één hashtabel in te voegen:
Gebruik ratio
Invoegduur 4 elementen
0,00
11,608448 ms (361,314798 miljoen toetsen/sec.)
0,03
11,751424 ms (356,918799 miljoen toetsen/sec.)
0,06
11,942592 ms (351,205515 miljoen toetsen/sec.)
0,09
12,081120 ms (347,178429 miljoen toetsen/sec.)
0,12
12,242560 ms (342,600233 miljoen toetsen/sec.)
0,16
12,396448 ms (338,347235 miljoen toetsen/sec.)
0,19
12,533024 ms (334,660176 miljoen toetsen/sec.)
0,22
12,703328 ms (330,173626 miljoen toetsen/sec.)
0,25
12,884512 ms (325,530693 miljoen toetsen/sec.)
0,28
13,033472 ms (321,810182 miljoen toetsen/sec.)
0,31
13,239296 ms (316,807174 miljoen toetsen/sec.)
0,34
13,392448 ms (313,184256 miljoen toetsen/sec.)
0,37
13,624000 ms (307,861434 miljoen toetsen/sec.)
0,41
13,875520 ms (302,280855 miljoen toetsen/sec.)
0,44
14,126528 ms (296,909756 miljoen toetsen/sec.)
0,47
14,399328 ms (291,284699 miljoen toetsen/sec.)
0,50
14,690304 ms (285,515123 miljoen toetsen/sec.)
0,53
15,039136 ms (278,892623 miljoen toetsen/sec.)
0,56
15,478656 ms (270,973402 miljoen toetsen/sec.)
0,59
15,985664 ms (262,379092 miljoen toetsen/sec.)
0,62
16,668673 ms (251,627968 miljoen toetsen/sec.)
0,66
17,587200 ms (238,486174 miljoen toetsen/sec.)
0,69
18,690048 ms (224,413765 miljoen toetsen/sec.)
0,72
20,278816 ms (206,831789 miljoen toetsen/sec.)
0,75
22,545408 ms (186,038058 miljoen toetsen/sec.)
0,78
26,053312 ms (160,989275 miljoen toetsen/sec.)
0,81
31,895008 ms (131,503463 miljoen toetsen/sec.)
0,84
42,103294 ms (99,619378 miljoen toetsen/sec.)
0,87
61,849056 ms (67,815164 miljoen toetsen/sec.)
0,90
105,695999 ms (39,682713 miljoen toetsen/sec.)
0,94
240,204636 ms (17,461378 miljoen toetsen/sec.)
Naarmate het gebruik toeneemt, nemen de prestaties af. Dit is in de meeste gevallen niet wenselijk. Als een applicatie elementen in een tabel invoegt en deze vervolgens weggooit (bijvoorbeeld bij het tellen van woorden in een boek), dan is dit geen probleem. Maar als de toepassing een hashtabel met een lange levensduur gebruikt (bijvoorbeeld in een grafische editor om niet-lege delen van afbeeldingen op te slaan waar de gebruiker regelmatig informatie invoegt en verwijdert), kan dit gedrag problematisch zijn.
En de sondediepte van de hashtabel gemeten na 64 miljoen inserts (gebruiksfactor 0,5). De gemiddelde diepte was 0,4774, dus de meeste toetsen bevonden zich in het best mogelijke slot of één slot verwijderd van de beste positie. De maximale peildiepte was 60.
Vervolgens heb ik de tastdiepte gemeten op een tafel met 124 miljoen inzetstukken (gebruiksfactor 0,97). De gemiddelde diepte was al 10,1757, en de maximale - 6474 (!!). De lineaire detectieprestaties nemen aanzienlijk af bij hoge bezettingsgraden.
Het is het beste om de bezettingsgraad van deze hashtabel laag te houden. Maar dan verhogen we de prestaties ten koste van het geheugengebruik. Gelukkig kan dit in het geval van 32-bits sleutels en waarden gerechtvaardigd worden. Als we in het bovenstaande voorbeeld in een tabel met 128 miljoen elementen de gebruiksfactor van 0,25 behouden, dan kunnen we er niet meer dan 32 miljoen elementen in plaatsen en gaan de resterende 96 miljoen slots verloren - 8 bytes voor elk paar , 768 MB verloren geheugen.
Houd er rekening mee dat we het hebben over het verlies van videokaartgeheugen, wat een waardevollere hulpbron is dan systeemgeheugen. Hoewel de meeste moderne desktop grafische kaarten die CUDA ondersteunen minimaal 4 GB geheugen hebben (op het moment van schrijven heeft de NVIDIA 2080 Ti 11 GB), zou het nog steeds niet de verstandigste beslissing zijn om dergelijke hoeveelheden te verliezen.
Later zal ik meer schrijven over het maken van hashtabellen voor videokaarten die geen problemen hebben met het onderzoeken van de diepte, en over manieren om verwijderde slots opnieuw te gebruiken.
Klinkende dieptemeting
Om de diepte van een sleutel te bepalen, kunnen we de hash van de sleutel (de ideale tabelindex) uit de werkelijke tabelindex halen:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Vanwege de magie van twee-twee-complement binaire getallen en het feit dat de capaciteit van de hashtabel twee tot de macht twee is, zal deze aanpak zelfs werken als de sleutelindex naar het begin van de tabel wordt verplaatst. Laten we een sleutel nemen die naar 1 is gehasht, maar in slot 3 is geplaatst. Dan krijgen we voor een tafel met capaciteit 4 (3 — 1) & 3, wat overeenkomt met 2.
Conclusie
Als u vragen of opmerkingen heeft, kunt u mij mailen op of open een nieuw onderwerp in.
Deze code is geschreven onder inspiratie van uitstekende artikelen:
In de toekomst zal ik blijven schrijven over hashtabelimplementaties voor videokaarten en hun prestaties analyseren. Mijn plannen omvatten chaining, Robin Hood-hashing en koekoek-hashing met behulp van atomaire operaties in datastructuren die GPU-vriendelijk zijn.
Bron: www.habr.com