Ich habe es auf Github gepostet .
Es handelt sich um eine einfache GPU-Hash-Tabelle, die Hunderte Millionen Einfügungen pro Sekunde verarbeiten kann. Auf meinem NVIDIA GTX 1060-Laptop fügt der Code in etwa 64 ms 210 Millionen zufällig generierte Schlüssel-Wert-Paare ein und entfernt 32 Millionen Paare in etwa 64 ms.
Das heißt, die Geschwindigkeit auf einem Laptop beträgt etwa 300 Millionen Einfügungen/Sekunde und 500 Millionen Löschungen/Sekunde.
Die Tabelle ist in CUDA geschrieben, obwohl die gleiche Technik auch auf HLSL oder GLSL angewendet werden kann. Die Implementierung unterliegt mehreren Einschränkungen, um eine hohe Leistung auf einer Grafikkarte sicherzustellen:
- Es werden nur 32-Bit-Schlüssel und die gleichen Werte verarbeitet.
- Die Hash-Tabelle hat eine feste Größe.
- Und diese Größe muss gleich zwei hoch sein.
Für Schlüssel und Werte müssen Sie eine einfache Trennmarkierung reservieren (im obigen Code ist dies 0xffffffff).
Hash-Tabelle ohne Sperren
Die Hash-Tabelle verwendet eine offene Adressierung mit Das heißt, es handelt sich einfach um ein Array von Schlüssel-Wert-Paaren, das im Speicher gespeichert wird und eine überlegene Cache-Leistung aufweist. Das Gleiche gilt nicht für die Verkettung, bei der nach einem Zeiger in einer verknüpften Liste gesucht wird. Eine Hash-Tabelle ist ein einfaches Array, in dem Elemente gespeichert sind KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Die Größe der Tabelle entspricht einer Zweierpotenz und nicht einer Primzahl, da ein schneller Befehl ausreicht, um die pow2/AND-Maske anzuwenden, der Moduloperator jedoch viel langsamer ist. Dies ist im Fall einer linearen Prüfung wichtig, da bei einer linearen Tabellensuche der Slot-Index in jeden Slot eingeschlossen werden muss. Infolgedessen addieren sich die Kosten des Vorgangs modulo in jedem Slot.
Die Tabelle speichert nur den Schlüssel und den Wert für jedes Element, keinen Hash des Schlüssels. Da die Tabelle nur 32-Bit-Schlüssel speichert, wird der Hash sehr schnell berechnet. Der obige Code verwendet den Murmur3-Hash, der nur wenige Verschiebungen, XORs und Multiplikationen durchführt.
Die Hash-Tabelle verwendet Sperrschutztechniken, die unabhängig von der Speicherreihenfolge sind. Selbst wenn einige Schreibvorgänge die Reihenfolge anderer solcher Vorgänge stören, behält die Hash-Tabelle dennoch den korrekten Zustand bei. Wir werden weiter unten darüber sprechen. Die Technik funktioniert hervorragend mit Grafikkarten, die Tausende von Threads gleichzeitig ausführen.
Die Schlüssel und Werte in der Hash-Tabelle werden leer initialisiert.
Der Code kann geändert werden, um auch 64-Bit-Schlüssel und -Werte zu verarbeiten. Schlüssel erfordern atomare Lese-, Schreib- und Vergleichs- und Austauschvorgänge. Und Werte erfordern atomare Lese- und Schreibvorgänge. Glücklicherweise sind in CUDA Lese-/Schreibvorgänge für 32- und 64-Bit-Werte atomar, solange sie auf natürliche Weise ausgerichtet sind (siehe unten). ) und moderne Grafikkarten unterstützen atomare 64-Bit-Vergleichs- und Austauschoperationen. Bei der Umstellung auf 64 Bit nimmt die Leistung natürlich leicht ab.
Status der Hash-Tabelle
Jedes Schlüssel-Wert-Paar in einer Hash-Tabelle kann einen von vier Zuständen haben:
- Schlüssel und Wert sind leer. In diesem Zustand wird die Hash-Tabelle initialisiert.
- Der Schlüssel wurde aufgeschrieben, der Wert wurde jedoch noch nicht geschrieben. Wenn ein anderer Thread gerade Daten liest, gibt er leer zurück. Das ist normal, das Gleiche wäre passiert, wenn ein anderer Ausführungsthread etwas früher funktioniert hätte, und wir sprechen von einer gleichzeitigen Datenstruktur.
- Sowohl der Schlüssel als auch der Wert werden aufgezeichnet.
- Der Wert ist für andere Ausführungsthreads verfügbar, der Schlüssel jedoch noch nicht. Dies kann passieren, weil das CUDA-Programmiermodell über ein lose geordnetes Speichermodell verfügt. Das ist normal; in jedem Fall ist der Schlüssel noch leer, auch wenn der Wert nicht mehr vorhanden ist.
Eine wichtige Nuance besteht darin, dass sich der Schlüssel, sobald er in den Steckplatz geschrieben wurde, nicht mehr bewegt – selbst wenn der Schlüssel gelöscht wird, werden wir weiter unten darüber sprechen.
Der Hash-Tabellencode funktioniert sogar mit lose geordneten Speichermodellen, bei denen die Reihenfolge, in der Speicher gelesen und geschrieben wird, unbekannt ist. Wenn wir uns das Einfügen, Suchen und Löschen in einer Hash-Tabelle ansehen, denken Sie daran, dass sich jedes Schlüssel-Wert-Paar in einem der vier oben beschriebenen Zustände befindet.
Einfügen in eine Hash-Tabelle
Die CUDA-Funktion, die Schlüssel-Wert-Paare in eine Hash-Tabelle einfügt, sieht folgendermaßen aus:
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);
}
}
Um einen Schlüssel einzufügen, durchläuft der Code das Hash-Tabellenarray, beginnend mit dem Hash des eingefügten Schlüssels. Jeder Slot im Array führt eine atomare Vergleichs- und Austauschoperation durch, die den Schlüssel in diesem Slot mit leer vergleicht. Wenn eine Nichtübereinstimmung festgestellt wird, wird der Schlüssel im Steckplatz mit dem eingesteckten Schlüssel aktualisiert und anschließend der ursprüngliche Steckplatzschlüssel zurückgegeben. Wenn dieser Originalschlüssel leer war oder mit dem eingefügten Schlüssel übereinstimmte, fand der Code einen geeigneten Steckplatz zum Einfügen und fügte den eingefügten Wert in den Steckplatz ein.
Wenn in einem Kernel-Aufruf gpu_hashtable_insert() Gibt es mehrere Elemente mit demselben Schlüssel, kann jeder ihrer Werte in den Schlüsselschlitz geschrieben werden. Dies gilt als normal: Einer der Schlüsselwert-Schreibvorgänge während des Aufrufs wird erfolgreich sein, aber da dies alles parallel in mehreren Ausführungsthreads geschieht, können wir nicht vorhersagen, welcher Speicherschreibvorgang der letzte sein wird.
Hash-Tabellensuche
Code für die Suche nach Schlüsseln:
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);
}
}
Um den Wert eines in einer Tabelle gespeicherten Schlüssels zu finden, durchlaufen wir das Array, beginnend mit dem Hash des gesuchten Schlüssels. In jedem Slot prüfen wir, ob der Schlüssel der gesuchte ist, und wenn ja, geben wir seinen Wert zurück. Wir prüfen auch, ob der Schlüssel leer ist, und brechen in diesem Fall die Suche ab.
Wenn wir den Schlüssel nicht finden können, gibt der Code einen leeren Wert zurück.
Alle diese Suchvorgänge können gleichzeitig durch Einfügungen und Löschungen ausgeführt werden. Jedes Paar in der Tabelle hat einen der vier oben für den Fluss beschriebenen Zustände.
Löschen in einer Hash-Tabelle
Code zum Löschen von Schlüsseln:
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);
}
}
Das Löschen eines Schlüssels erfolgt auf ungewöhnliche Weise: Wir belassen den Schlüssel in der Tabelle und markieren seinen Wert (nicht den Schlüssel selbst) als leer. Dieser Code ist sehr ähnlich lookup(), mit der Ausnahme, dass, wenn eine Übereinstimmung für einen Schlüssel gefunden wird, dessen Wert leer wird.
Wie oben erwähnt, wird ein Schlüssel, sobald er in einen Steckplatz geschrieben wurde, nicht mehr verschoben. Selbst wenn ein Element aus der Tabelle gelöscht wird, bleibt der Schlüssel bestehen, sein Wert wird einfach leer. Dies bedeutet, dass wir keine atomare Schreiboperation für den Slot-Wert verwenden müssen, da es keine Rolle spielt, ob der aktuelle Wert leer ist oder nicht, er wird trotzdem leer.
Größe einer Hash-Tabelle ändern
Sie können die Größe einer Hash-Tabelle ändern, indem Sie eine größere Tabelle erstellen und nicht leere Elemente aus der alten Tabelle in diese einfügen. Ich habe diese Funktionalität nicht implementiert, weil ich den Beispielcode einfach halten wollte. Darüber hinaus erfolgt die Speicherzuweisung in CUDA-Programmen häufig im Hostcode und nicht im CUDA-Kernel.
Artikel beschreibt, wie eine solche sperrengeschützte Datenstruktur geändert wird.
Wettbewerbsfähigkeit
In den obigen Funktionscodeausschnitten gpu_hashtable_insert(), _lookup() и _delete() Verarbeiten Sie jeweils ein Schlüssel-Wert-Paar. Und niedriger gpu_hashtable_insert(), _lookup() и _delete() Verarbeiten Sie ein Array von Paaren parallel, jedes Paar in einem separaten GPU-Ausführungsthread:
// 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);
}
}
Die sperrenresistente Hash-Tabelle unterstützt gleichzeitige Einfügungen, Suchvorgänge und Löschungen. Da sich Schlüssel-Wert-Paare immer in einem von vier Zuständen befinden und die Schlüssel sich nicht bewegen, garantiert die Tabelle die Korrektheit auch dann, wenn verschiedene Arten von Operationen gleichzeitig verwendet werden.
Wenn wir jedoch eine Reihe von Einfügungen und Löschungen parallel verarbeiten und das Eingabearray von Paaren doppelte Schlüssel enthält, können wir nicht vorhersagen, welche Paare „gewinnen“ werden – sie werden zuletzt in die Hash-Tabelle geschrieben. Nehmen wir an, wir haben den Einfügecode mit einem Eingabearray aus Paaren aufgerufen A/0 B/1 A/2 C/3 A/4. Wenn der Code abgeschlossen ist, erfolgt die Paarung B/1 и C/3 sind garantiert in der Tabelle vorhanden, aber gleichzeitig erscheint jedes der Paare darin A/0, A/2 oder A/4. Dies kann ein Problem sein oder auch nicht – es hängt alles von der Anwendung ab. Möglicherweise wissen Sie im Voraus, dass das Eingabearray keine doppelten Schlüssel enthält, oder es ist Ihnen möglicherweise egal, welcher Wert zuletzt geschrieben wurde.
Wenn dies für Sie ein Problem darstellt, müssen Sie die doppelten Paare in verschiedene CUDA-Systemaufrufe aufteilen. In CUDA wird jede Operation, die den Kernel aufruft, immer vor dem nächsten Kernel-Aufruf abgeschlossen (zumindest innerhalb eines Threads. In verschiedenen Threads werden Kernel parallel ausgeführt). Wenn Sie im obigen Beispiel einen Kernel mit aufrufen A/0 B/1 A/2 C/3, und der andere mit A/4, dann der Schlüssel A wird den Wert bekommen 4.
Lassen Sie uns nun darüber sprechen, ob Funktionen dies tun sollten lookup() и delete() Verwenden Sie einen einfachen oder flüchtigen Zeiger auf ein Array von Paaren in der Hash-Tabelle. Besagt, dass:
Der Compiler kann sich dafür entscheiden, Lese- und Schreibvorgänge im globalen oder gemeinsam genutzten Speicher zu optimieren. Diese Optimierungen können mit dem Schlüsselwort deaktiviert werden
volatile: ... jeder Verweis auf diese Variable wird in eine echte Speicher-Lese- oder Schreibanweisung kompiliert.
Korrektheitserwägungen bedürfen keiner Anwendung volatile. Wenn der Ausführungsthread einen zwischengespeicherten Wert aus einem früheren Lesevorgang verwendet, verwendet er leicht veraltete Informationen. Dennoch handelt es sich hierbei um Informationen aus dem korrekten Zustand der Hash-Tabelle zu einem bestimmten Zeitpunkt des Kernel-Aufrufs. Wenn Sie die neuesten Informationen benötigen, können Sie den Index verwenden volatile, aber dann nimmt die Leistung leicht ab: Laut meinen Tests sank die Geschwindigkeit beim Löschen von 32 Millionen Elementen von 500 Millionen Löschungen/Sek. auf 450 Millionen Löschungen/Sek.
Leistung
Im Test zum Einfügen von 64 Millionen Elementen und zum Löschen von 32 Millionen davon herrscht Konkurrenz zwischen std::unordered_map und es gibt praktisch keine Hash-Tabelle für die GPU:
std::unordered_map verbrachte 70 ms damit, Elemente einzufügen, zu entfernen und sie dann freizugeben unordered_map (Millionen Elemente loszuwerden, kostet viel Zeit, denn im Inneren unordered_map es werden mehrere Speicherzuweisungen vorgenommen). Ehrlich gesagt, std:unordered_map ganz andere Einschränkungen. Es handelt sich um einen einzelnen CPU-Ausführungsthread, der Schlüsselwerte jeder Größe unterstützt, bei hohen Auslastungsraten eine gute Leistung erbringt und nach mehreren Löschvorgängen eine stabile Leistung zeigt.
Die Dauer der Hash-Tabelle für die GPU- und Interprogramm-Kommunikation betrug 984 ms. Dazu gehört die Zeit, die für das Platzieren und Löschen der Tabelle im Speicher (einmaliges Zuweisen von 1 GB Speicher, was in CUDA einige Zeit dauert), für das Einfügen und Löschen von Elementen sowie für das Durchlaufen dieser Elemente aufgewendet wird. Alle Kopien zum und vom Grafikkartenspeicher werden ebenfalls berücksichtigt.
Die Fertigstellung der Hash-Tabelle selbst dauerte 271 ms. Dazu gehört auch die Zeit, die die Grafikkarte für das Einfügen und Löschen von Elementen benötigt, und berücksichtigt nicht die Zeit, die für das Kopieren in den Speicher und das Durchlaufen der resultierenden Tabelle aufgewendet wird. Wenn die GPU-Tabelle lange existiert oder wenn die Hash-Tabelle vollständig im Speicher der Grafikkarte enthalten ist (z. B. um eine Hash-Tabelle zu erstellen, die von anderem GPU-Code und nicht vom Zentralprozessor verwendet wird), dann das Testergebnis ist relevant.
Die Hash-Tabelle für eine Grafikkarte weist aufgrund des hohen Durchsatzes und der aktiven Parallelisierung eine hohe Leistung auf.
Begrenztheit
Bei der Hash-Tabellenarchitektur sind einige Probleme zu beachten:
- Die lineare Prüfung wird durch Clustering behindert, was dazu führt, dass die Schlüssel in der Tabelle nicht perfekt platziert sind.
- Schlüssel werden mit der Funktion nicht entfernt
deleteund mit der Zeit verstopfen sie den Tisch.
Infolgedessen kann sich die Leistung einer Hash-Tabelle allmählich verschlechtern, insbesondere wenn sie über einen längeren Zeitraum existiert und zahlreiche Einfügungen und Löschungen aufweist. Eine Möglichkeit, diese Nachteile abzumildern, besteht darin, eine neue Tabelle mit einer relativ geringen Auslastungsrate erneut aufzubereiten und die entfernten Schlüssel während der erneuten Aufbereitung herauszufiltern.
Um die beschriebenen Probleme zu veranschaulichen, verwende ich den obigen Code, um eine Tabelle mit 128 Millionen Elementen zu erstellen und 4 Millionen Elemente zu durchlaufen, bis ich 124 Millionen Slots gefüllt habe (Auslastungsrate von etwa 0,96). Hier ist eine Ergebnistabelle. Jede Zeile ist ein Aufruf an den CUDA-Kernel, 4 Millionen neue Elemente in eine Hash-Tabelle einzufügen:
Benutzungsrate
Einfügungsdauer 4 Elemente
0,00
11,608448 ms (361,314798 Millionen Schlüssel/Sek.)
0,03
11,751424 ms (356,918799 Millionen Schlüssel/Sek.)
0,06
11,942592 ms (351,205515 Millionen Schlüssel/Sek.)
0,09
12,081120 ms (347,178429 Millionen Schlüssel/Sek.)
0,12
12,242560 ms (342,600233 Millionen Schlüssel/Sek.)
0,16
12,396448 ms (338,347235 Millionen Schlüssel/Sek.)
0,19
12,533024 ms (334,660176 Millionen Schlüssel/Sek.)
0,22
12,703328 ms (330,173626 Millionen Schlüssel/Sek.)
0,25
12,884512 ms (325,530693 Millionen Schlüssel/Sek.)
0,28
13,033472 ms (321,810182 Millionen Schlüssel/Sek.)
0,31
13,239296 ms (316,807174 Millionen Schlüssel/Sek.)
0,34
13,392448 ms (313,184256 Millionen Schlüssel/Sek.)
0,37
13,624000 ms (307,861434 Millionen Schlüssel/Sek.)
0,41
13,875520 ms (302,280855 Millionen Schlüssel/Sek.)
0,44
14,126528 ms (296,909756 Millionen Schlüssel/Sek.)
0,47
14,399328 ms (291,284699 Millionen Schlüssel/Sek.)
0,50
14,690304 ms (285,515123 Millionen Schlüssel/Sek.)
0,53
15,039136 ms (278,892623 Millionen Schlüssel/Sek.)
0,56
15,478656 ms (270,973402 Millionen Schlüssel/Sek.)
0,59
15,985664 ms (262,379092 Millionen Schlüssel/Sek.)
0,62
16,668673 ms (251,627968 Millionen Schlüssel/Sek.)
0,66
17,587200 ms (238,486174 Millionen Schlüssel/Sek.)
0,69
18,690048 ms (224,413765 Millionen Schlüssel/Sek.)
0,72
20,278816 ms (206,831789 Millionen Schlüssel/Sek.)
0,75
22,545408 ms (186,038058 Millionen Schlüssel/Sek.)
0,78
26,053312 ms (160,989275 Millionen Schlüssel/Sek.)
0,81
31,895008 ms (131,503463 Millionen Schlüssel/Sek.)
0,84
42,103294 ms (99,619378 Millionen Schlüssel/Sek.)
0,87
61,849056 ms (67,815164 Millionen Schlüssel/Sek.)
0,90
105,695999 ms (39,682713 Millionen Schlüssel/Sek.)
0,94
240,204636 ms (17,461378 Millionen Schlüssel/Sek.)
Mit zunehmender Auslastung nimmt die Leistung ab. Dies ist in den meisten Fällen nicht erwünscht. Wenn eine Anwendung Elemente in eine Tabelle einfügt und diese dann verwirft (z. B. beim Zählen von Wörtern in einem Buch), ist dies kein Problem. Wenn die Anwendung jedoch eine langlebige Hash-Tabelle verwendet (z. B. in einem Grafikeditor zum Speichern nicht leerer Teile von Bildern, in denen der Benutzer häufig Informationen einfügt und löscht), kann dieses Verhalten problematisch sein.
Und die Prüftiefe der Hash-Tabelle wurde nach 64 Millionen Einfügungen gemessen (Nutzungsfaktor 0,5). Die durchschnittliche Tiefe betrug 0,4774, sodass sich die meisten Tasten entweder im bestmöglichen Steckplatz oder einen Steckplatz von der besten Position entfernt befanden. Die maximale Sondierungstiefe betrug 60.
Anschließend habe ich die Antasttiefe an einem Tisch mit 124 Millionen Einsätzen gemessen (Nutzungsfaktor 0,97). Die durchschnittliche Tiefe betrug bereits 10,1757 und die maximale - 6474 (!!). Bei hohen Auslastungsraten sinkt die Leistung der linearen Sensorik erheblich.
Es ist am besten, die Auslastungsrate dieser Hash-Tabelle niedrig zu halten. Aber dann steigern wir die Leistung auf Kosten des Speicherverbrauchs. Glücklicherweise kann dies bei 32-Bit-Schlüsseln und -Werten gerechtfertigt sein. Wenn wir im obigen Beispiel in einer Tabelle mit 128 Millionen Elementen den Auslastungsfaktor von 0,25 beibehalten, können wir nicht mehr als 32 Millionen Elemente darin platzieren und die verbleibenden 96 Millionen Slots gehen verloren – 8 Bytes für jedes Paar , 768 MB verlorener Speicher.
Bitte beachten Sie, dass es sich um den Verlust des Grafikkartenspeichers handelt, der eine wertvollere Ressource darstellt als der Systemspeicher. Obwohl die meisten modernen Desktop-Grafikkarten, die CUDA unterstützen, über mindestens 4 GB Speicher verfügen (zum Zeitpunkt des Schreibens verfügt die NVIDIA 2080 Ti über 11 GB), wäre es dennoch nicht die klügste Entscheidung, solche Mengen zu verlieren.
Später werde ich mehr über das Erstellen von Hash-Tabellen für Grafikkarten schreiben, bei denen es keine Probleme mit der Prüftiefe gibt, sowie über Möglichkeiten zur Wiederverwendung gelöschter Steckplätze.
Sondierungstiefenmessung
Um die Prüftiefe eines Schlüssels zu bestimmen, können wir den Hash des Schlüssels (seinen idealen Tabellenindex) aus seinem tatsächlichen Tabellenindex extrahieren:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Aufgrund der Magie der Zweierkomplement-Binärzahlen und der Tatsache, dass die Kapazität der Hash-Tabelle zwei hoch zwei beträgt, funktioniert dieser Ansatz auch dann, wenn der Schlüsselindex an den Anfang der Tabelle verschoben wird. Nehmen wir einen Schlüssel, der auf 1 gehasht ist, aber in Steckplatz 3 eingefügt wird. Dann erhalten wir für eine Tabelle mit der Kapazität 4 (3 — 1) & 3, was 2 entspricht.
Abschluss
Wenn Sie Fragen oder Kommentare haben, senden Sie mir bitte eine E-Mail an oder eröffnen Sie ein neues Thema in.
Dieser Code wurde unter Inspiration von hervorragenden Artikeln geschrieben:
In Zukunft werde ich weiterhin über Hash-Tabellen-Implementierungen für Grafikkarten schreiben und deren Leistung analysieren. Zu meinen Plänen gehören Verkettung, Robin-Hood-Hashing und Cuckoo-Hashing mit atomaren Operationen in Datenstrukturen, die GPU-freundlich sind.
Source: habr.com