Ich habe es auf Github veröffentlicht .
Dies ist eine einfache Hash-Tabelle für GPU, die in der Lage ist, Hundert Millionen Einfügungen pro Sekunde zu verarbeiten. Auf meinem Laptop mit NVIDIA GTX 1060 fügt der Code etwa 64 Millionen zufällig generierter Schlüssel-Wert-Paare in ungefähr 210 ms ein und entfernt 32 Millionen Paare in etwa 64 ms.
Das heißt, die Geschwindigkeit auf dem Laptop beträgt etwa 300 Millionen Einfügungen/Sekunde und 500 Millionen Löschungen/Sekunde.
Die Tabelle ist in CUDA geschrieben, obwohl dieselbe Methode auf HLSL oder GLSL angewendet werden kann. Die Implementierung hat einige Einschränkungen, die eine hohe Leistung auf der Grafikkarte gewährleisten:
- Es werden nur 32-Bit-Schlüssel und die entsprechenden Werte verarbeitet.
- Die Hash-Tabelle hat eine feste Größe.
- Und diese Größe muss eine Potenz von zwei sein.
Für Schlüssel und Werte muss ein einfacher Trennmarker reserviert werden (im bereitgestellten Code ist dies 0xffffffff).
Lockfreie Hash-Tabelle
In der Hash-Tabelle wird offene Adressierung mit verwendet, d. h. es handelt sich einfach um ein Array von Schlüssel-Wert-Paaren, das im Speicher gespeichert ist und eine hervorragende Cache-Leistung aufweist. Das lässt sich nicht über das Verkettungssystem (chaining) sagen, bei dem ein Zeiger in einer verketteten Liste gesucht werden muss. Die Hash-Tabelle ist ein einfaches Array, das Elemente speichert. KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Die Größe der Tabelle ist eine Potenz von zwei und nicht eine Primzahl, weil für die Anwendung von pow2/AND-Maske nur eine schnelle Anweisung erforderlich ist, während der Modulo-Operator viel langsamer arbeitet. Dies ist wichtig beim linearen Sondieren, da beim linearen Suchen durch die Tabelle der Slot-Index auf jeden Slot gewickelt werden muss. Und infolgedessen wird die Kosten für die Modulo-Operation in jedem Slot hinzugefügt.
Die Tabelle speichert nur den Schlüssel und den Wert für jedes Element, nicht den Hash des Schlüssels. Da die Tabelle nur 32-Bit-Schlüssel speichert, wird der Hash sehr schnell berechnet. Im bereitgestellten Code wird der Hash Murmur3 verwendet, der nur einige Verschiebungen, XOR-Operationen und Multiplikationen durchführt.
In der Hash-Tabelle wird eine Schutztechnik gegen Blockaden angewandt, die unabhängig von der Anordnung im Speicher ist. Selbst wenn einige Schreiboperationen die Reihenfolge anderer solcher Operationen stören, bleibt der Zustand der Hash-Tabelle dennoch korrekt. Darüber werden wir später sprechen. Die Technik funktioniert hervorragend mit Grafikkarten, in denen Tausende von Threads gleichzeitig arbeiten.
Die Schlüssel und Werte in der Hash-Tabelle werden als leer initialisiert.
Der Code kann modifiziert werden, um sowohl 64-Bit-Schlüssel als auch -Werte zu verarbeiten. Für die Schlüssel sind atomare Lese-, Schreib- und Vergleichsoperationen mit Austausch (compare-and-swap) erforderlich. Für die Werte sind atomare Lese- und Schreiboperationen notwendig. Glücklicherweise sind in CUDA die Lese-Schreib-Operationen für 32- und 64-Bit-Werte atomar, solange sie natürlich ausgerichtet sind (siehe ), und moderne Grafikkarten unterstützen 64-Bit-atomare Vergleichsoperationen mit Austausch. Natürlich wird die Leistung beim Wechsel auf 64 Bit etwas sinken.
Zustand der Hash-Tabelle
Jedes Schlüssel-Wert-Paar in der 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 geschrieben, der Wert jedoch noch nicht. Wenn zu diesem Zeitpunkt ein anderer Ausführungsthread die Daten liest, wird er einen leeren Wert zurückgeben. Das ist in Ordnung, dasselbe würde passieren, wenn ein anderer Ausführungsthread etwas früher ausgeführt wird, und wir sprechen über eine konkurrierende Datenstruktur.
- Sowohl der Schlüssel als auch der Wert sind geschrieben.
- Der Wert ist für andere Ausführungsthreads verfügbar, der Schlüssel jedoch noch nicht. Dies kann vorkommen, weil das Programmiermodell in CUDA ein schwach geordnetes Modell des Speichers unterstellt. Das ist in Ordnung, unter allen Umständen bleibt der Schlüssel immer noch leer, selbst wenn der Wert bereits nicht mehr so ist.
Ein wichtiger Punkt ist, dass, sobald der Schlüssel in den Slot geschrieben wurde, er nicht mehr verschoben wird – selbst wenn der Schlüssel gelöscht wird, darüber werden wir später sprechen.
Der Code für die Hash-Tabelle funktioniert sogar mit schwach geordneten Modellen des Speichers, bei denen die Reihenfolge von Lese- und Schreibvorgängen im Speicher nicht bekannt ist. Wenn wir die Einfügung, Suche und Löschung in der Hash-Tabelle durchgehen, denken Sie daran, dass jedes Schlüssel-Wert-Paar in einem der vier oben beschriebenen Zustände ist.
Einfügen in die 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);
}
}
Für die Einfügung eines Schlüssels iteriert der Code durch das Array der Hash-Tabelle, beginnend mit dem Hash des einzufügenden Schlüssels. In jedem Slot des Arrays wird eine atomare Vergleichs-und-Tausch-Operation durchgeführt, bei der der Schlüssel in diesem Slot mit leer verglichen wird. Wird eine Abweichung festgestellt, wird der Schlüssel im Slot durch den einzufügenden Schlüssel aktualisiert, und dann wird der ursprüngliche Schlüssel des Slots zurückgegeben. War dieser ursprüngliche Schlüssel leer oder entsprach er dem einzufügenden Schlüssel, hat der Code den passenden Slot zum Einfügen gefunden und fügt den Wert in den Slot ein.
Wenn in einem Kernaufruf gpu_hashtable_insert() mehrere Elemente mit demselben Schlüssel vorhanden sind, kann eines ihrer Werte in den Schlüsselslot geschrieben werden. Dies wird als normal angesehen: Eine der Schlüssel-Wert-Schreiboperationen während des Aufrufs wird erfolgreich sein, aber da dies alles parallel in mehreren Ausführungsthreads geschieht, können wir nicht vorhersagen, welche Speicheroperation zuletzt ausgeführt wird.
Suche in der Hash-Tabelle
Code zum Suchen von 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 Schlüssels zu finden, der in der Tabelle gespeichert ist, iterieren wir durch 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. Außerdem prüfen wir, ob der Schlüssel leer ist, und wenn ja, brechen wir die Suche ab.
Wenn wir den Schlüssel nicht finden können, gibt der Code einen leeren Wert zurück.
All diese Suchoperationen können gleichzeitig während Einfügungen und Löschungen ausgeführt werden. Jedes Paar in der Tabelle hat für den Thread einen der vier oben beschriebenen Zustände.
Löschen in der 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 lassen den Schlüssel in der Tabelle und markieren seinen Wert (nicht den Schlüssel selbst) als leer. Dieser Code ähnelt sehr lookup(), mit der Ausnahme, dass beim Finden einer Übereinstimmung mit dem Schlüssel dessen Wert leer wird.
Wie bereits erwähnt, sobald der Schlüssel in den Slot geschrieben wurde, wird er nicht mehr verschoben. Selbst beim Löschen eines Elements aus der Tabelle bleibt der Schlüssel an seinem Platz, sein Wert wird einfach leer. Das bedeutet, dass wir keine atomare Operation zum Schreiben des Slot-Werts verwenden müssen, da es egal ist, ob der aktuelle Wert leer ist oder nicht – er wird trotzdem leer.
Ändern der Größe der Hash-Tabelle
Die Größe der Hash-Tabelle kann durch Erstellen einer größeren Tabelle und Einfügen der nicht-leeren Elemente aus der alten Tabelle geändert werden. Ich habe diese Funktionalität nicht implementiert, da ich den Code einfach halten wollte. Darüber hinaus wird in CUDA-Programmen die Speicherzuweisung häufig im Host-Code und nicht im CUDA-Kernel durchgeführt.
In dem Artikel wird beschrieben, wie eine solche von Sperren geschützte Datenstruktur geändert werden kann.
Konkurrenz
In den oben genannten Codefragmenten verarbeiten die Funktionen gpu_hashtable_insert(), _lookup() und _delete() jeweils ein Schlüssel-Wert-Paar gleichzeitig. Unten werden gpu_hashtable_insert(), _lookup() und _delete() Paare von Elementen parallel verarbeitet, jedes Paar in einem separaten GPU-Thread:
// 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 Hash-Tabelle mit Sperrensicherung unterstützt konkurrierende Einfügungen, Suchanfragen und Löschvorgänge. Da sich die Schlüssel-Wert-Paare immer in einem von vier Zuständen befinden und die Schlüssel nicht verschoben werden, garantiert die Tabelle die Korrektheit selbst bei gleichzeitiger Verwendung von Operationen verschiedener Art.
Wenn wir jedoch gleichzeitig ein Paket von Einfügungen und Löschungen verarbeiten und wenn das Eingangsarray von Paaren doppelt vorhandene Schlüssel enthält, können wir nicht vorhersagen, welche Paare "gewinnen" – die zuletzt in die Hash-Tabelle geschrieben werden. Angenommen, wir haben den Einfügecode mit dem Eingangsarray von Paaren aufgerufen, A/0 B/1 A/2 C/3 A/4. Wenn der Code abgeschlossen ist, werden die Paare B/1 und C/3 mit Sicherheit in der Tabelle vorhanden sein, aber es wird auch eines der Paare A/0, A/2 oder A/4. Dies kann ein Problem sein oder auch nicht – es hängt alles von der Anwendung ab. Sie können vorab wissen, dass im Eingabearray keine doppelten Schlüssel vorhanden sind, oder es kann Ihnen egal sein, welches Wert zuletzt geschrieben wurde.
Wenn das für Sie ein Problem darstellt, müssen Sie doppelte Paare in verschiedene systemeigene CUDA-Aufrufe aufteilen. In CUDA wird jede Operation mit einem Kernelaufruf immer abgeschlossen, bevor der nächste Kernelaufruf erfolgt (zumindest innerhalb eines Threads. In verschiedenen Threads werden die Kerne parallel ausgeführt). Wenn im obigen Beispiel ein Kernel mit A/0 B/1 A/2 C/3, und ein anderer mit A/4, dann erhält der Schlüssel A den Wert 4.
Jetzt sprechen wir darüber, ob die Funktionen lookup() und delete() einen einfachen (plain) oder variablen (volatile) Zeiger auf das Array von Paaren in der Hash-Tabelle verwenden sollten. stellt fest, dass:
Der Compiler kann Lese- und Schreiboperationen im globalen oder gemeinsamen Speicher nach eigenem Ermessen optimieren … Diese Optimierungen können mit dem Schlüsselwort
volatile: … jeder Verweis auf diese Variable wird in eine tatsächliche Lese- oder Schreiboperation im Speicher kompiliert.
Die Überlegungen zur Korrektheit erfordern keinen Einsatz von volatile. Wenn der Ausführungsthread einen zwischengespeicherten Wert aus einer früheren Leseoperation verwendet, bedeutet das, dass er ein wenig veraltete Informationen nutzt. Dennoch handelt es sich um Informationen aus einem gültigen Zustand der Hash-Tabelle zu einem bestimmten Zeitpunkt des Kernelaufrufs. Wenn Sie die aktuellsten Informationen benötigen, können Sie einen Zeiger verwenden volatile, aber dann wird die Leistung etwas sinken: Nach meinen Tests verringerte sich beim Löschen von 32 Millionen Elementen die Geschwindigkeit von 500 Millionen Löschungen/Sekunde auf 450 Millionen Löschungen/Sekunde.
Leistung
Im Test mit dem Einfügen von 64 Millionen Elementen und dem Löschen von 32 Millionen daraus gab es keine Konkurrenz zwischen std::unordered_map und der Hash-Tabelle für die GPU:
std::unordered_map benötigte sie 70.691 ms für das Einfügen und Löschen von Elementen mit anschließender Freigabe von unordered_map (die Freigabe von Millionen von Elementen dauert einiges an Zeit, da zahlreiche Speicherzuweisungen innerhalb unordered_map stattfinden). Ehrlich gesagt, bei std:unordered_map ganz andere Einschränkungen. Es handelt sich um einen einzigen CPU-Ausführungs-Thread, der Schlüssel-Wert-Paare beliebiger Größe unterstützt, gut bei hohen Nutzungsraten funktioniert und nach zahlreichen Löschvorgängen eine stabile Leistung zeigt.
Die Laufzeit der Hash-Tabelle für GPU und Interprozesskommunikation betrug 984 ms. Darin enthalten sind die Zeit für die Platzierung der Tabelle im Speicher und deren Löschung (einmalige Zuweisung von 1 GB Speicher, was in CUDA einige Zeit in Anspruch nimmt), das Einfügen und Löschen von Elementen sowie das Iterieren über diese. Außerdem wurden alle Kopiervorgänge in den Speicher und aus dem Speicher der Grafikkarte berücksichtigt.
Die Verarbeitung der Hash-Tabelle selbst benötigte 271 ms. Es ist die Zeit, die die Grafikkarte für das Einfügen und Löschen von Elementen aufwendet, und die Zeit für das Kopieren in den Speicher und das Iterieren über die resultierende Tabelle ist nicht enthalten. Wenn die GPU-Tabelle lange lebt oder wenn die Hash-Tabelle vollständig im Speicher der Grafikkarte enthalten ist (zum Beispiel für die Erstellung einer Hash-Tabelle, die von anderem GPU-Code und nicht vom zentralen Prozessor verwendet wird), ist das Testergebnis relevant.
Die Hash-Tabelle für die Grafikkarte zeigt eine hohe Leistung dank der hohen Bandbreite und aktiven Parallelisierung.
Nachteile
Die Architektur der Hash-Tabelle weist mehrere Probleme auf, die berücksichtigt werden müssen:
- Lineares Abtasten wird durch Clusterbildung behindert, wodurch Schlüssel in der Tabelle alles andere als ideal platziert werden.
- Schlüssel werden nicht mit der Funktion
deleteund überladen mit der Zeit die Tabelle.
Infolgedessen kann die Leistung der Hash-Tabelle allmählich abnehmen, insbesondere wenn sie lange besteht und zahlreiche Einfügungen und Löschungen durchgeführt werden. Eine Möglichkeit zur Minderung dieser Nachteile besteht darin, in eine neue Tabelle mit einem ausreichend niedrigen Nutzungsgrad umzuhebeln und beim Umsetzen gelöschte Schlüssel herauszufiltern.
Um die beschriebenen Probleme zu veranschaulichen, verwende ich den oben genannten Code zur Erstellung einer Tabelle mit 128 Millionen Elementen, in die ich zyklisch 4 Millionen Elemente einfüge, bis ich 124 Millionen Slots gefüllt habe (Nutzungsgrad etwa 0,96). Hier ist die Ergebnistabelle, jede Zeile ist ein Aufruf des CUDA-Kernels mit dem Einfügen von 4 Millionen neuen Elementen in eine Hash-Tabelle:
Nutzungsgrad
Dauer der Einfügung von 4 194 304 Elementen
0,00
11,608448 ms (361,314798 Mio. Schlüssel/Sek.)
0,03
11,751424 ms (356,918799 Mio. Schlüssel/Sek.)
0,06
11,942592 ms (351,205515 Mio. Schlüssel/Sek.)
0,09
12,081120 ms (347,178429 Mio. Schlüssel/Sek.)
0,12
12,242560 ms (342,600233 Mio. Schlüssel/Sek.)
0,16
12,396448 ms (338,347235 Mio. Schlüssel/Sek.)
0,19
12,533024 ms (334,660176 Mio. Schlüssel/Sek.)
0,22
12,703328 ms (330,173626 Mio. Schlüssel/Sek.)
0,25
12,884512 ms (325,530693 Mio. Schlüssel/Sek.)
0,28
13,033472 ms (321,810182 Mio. Schlüssel/Sek.)
0,31
13,239296 ms (316,807174 Mio. Schlüssel/Sek.)
0,34
13,392448 ms (313,184256 Mio. Schlüssel/Sek.)
0,37
13,624000 ms (307,861434 Mio. Schlüssel/Sek.)
0,41
13,875520 ms (302,280855 Mio. Schlüssel/Sek.)
0,44
14,126528 ms (296,909756 Mio. Schlüssel/Sek.)
0,47
14,399328 ms (291,284699 Mio. Schlüssel/Sek.)
0,50
14,690304 ms (285,515123 Mio. Schlüssel/Sek.)
0,53
15,039136 ms (278,892623 Mio. Schlüssel/Sek.)
0,56
15,478656 ms (270,973402 Mio. Schlüssel/Sek.)
0,59
15,985664 ms (262,379092 Mio. Schlüssel/Sek.)
0,62
16,668673 ms (251,627968 Mio. Schlüssel/Sek.)
0,66
17,587200 ms (238,486174 Mio. Schlüssel/Sek.)
0,69
18,690048 ms (224,413765 Mio. Schlüssel/Sek.)
0,72
20,278816 ms (206,831789 Mio. Schlüssel/Sek.)
0,75
22,545408 ms (186,038058 Mio. Schlüssel/Sek.)
0,78
26,053312 ms (160,989275 Mio. Schlüssel/Sek.)
0,81
31,895008 ms (131,503463 Mio. Schlüssel/Sek.)
0,84
42,103294 ms (99,619378 Mio. Schlüssel/Sek.)
0,87
61,849056 ms (67,815164 Mio. Schlüssel/Sek.)
0,90
105,695999 ms (39,682713 Mio. Schlüssel/Sek.)
0,94
240,204636 ms (17,461378 Mio. Schlüssel/Sek.)
Mit steigendem Auslastungsgrad sinkt die Leistung. Dies ist in den meisten Fällen unerwünscht. Wenn eine Anwendung Elemente in einer Tabelle einfügt und sie dann verwirft (zum Beispiel beim Zählen von Wörtern in einem Buch), ist das kein Problem. Aber wenn die Anwendung eine langlebige Hash-Tabelle verwendet (zum Beispiel in einem Grafikeditor zur Speicherung nicht leerer Teile von Bildern, wenn der Benutzer häufig Informationen einfügt und löscht), kann dieses Verhalten störend sein.
Ich habe die Durchdringungstiefe der Hash-Tabelle nach 64 Millionen Einfügungen (Nutzung 0,5) gemessen. Die durchschnittliche Tiefe betrug 0,4774, sodass sich die meisten Schlüssel entweder in dem besten möglichen Slot oder in einem Slot vom besten Platz befanden. Die maximale Durchdringungstiefe betrug 60.
Dann habe ich die Durchdringungstiefe in einer Tabelle mit 124 Millionen Einfügungen (Nutzung 0,97) gemessen. Die durchschnittliche Tiefe betrug bereits 10,1757, und die maximale — 6474 (!!). Die Leistung der linearen Durchdringung nimmt bei hohen Auslastungsgraden stark ab.
Es ist am besten, bei dieser Hash-Tabelle einen niedrigen Auslastungsgrad beizubehalten. Aber dann erhöhen wir die Leistung auf Kosten des Speicherverbrauchs. Zum Glück kann dies im Fall von 32-Bit-Schlüsseln und -Werten gerechtfertigt sein. Wenn wir im obigen Beispiel bei einer Tabelle mit 128 Millionen Elementen einen Auslastungsgrad von 0,25 beibehalten, können wir nicht mehr als 32 Millionen Elemente unterbringen, während die restlichen 96 Millionen Slots verloren gehen - 8 Byte pro Paar, was 768 MB verlorenen Speicherplatz bedeutet.
Bitte beachten Sie, dass es hier um den Speicherverlust der Grafikkarte geht, die ein wertvoller Ressourcen ist als der Systemspeicher. Obwohl die meisten modernen Desktop-Grafikkarten, die CUDA unterstützen, mindestens 4 GB Speicher haben (zum Zeitpunkt des Schreibens hat NVIDIA 2080 Ti 11 GB), wäre es dennoch nicht weise, solche Mengen zu verlieren.
Später werde ich ausführlicher über die Erstellung von Hash-Tabellen für Grafikkarten schreiben, die keine Probleme mit der Tiefe der Sondierung haben, sowie über Möglichkeiten zur Wiederverwendung gelöschter Slots.
Sondierungstiefe messen
Um die Sondierungstiefe eines Schlüssels zu bestimmen, können wir den Hash des Schlüssels (seinen idealen Index in der Tabelle) von seinem tatsächlichen Tabellenindex abziehen:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Aufgrund der Magie zweier Binärzahlen im Zweierkomplement und der Tatsache, dass die Kapazität der Hash-Tabelle eine Potenz von zwei ist, funktioniert dieser Ansatz sogar dann, wenn der Schlüsselindex an den Anfang der Tabelle verschoben wird. Nehmen wir an, der Schlüssel wird auf 1 gehasht, aber in Slot 3 eingefügt. Bei einer Tabelle mit einer Kapazität von 4 erhalten wir dann (3 - 1) & 3, was 2 entspricht.
Fazit
Wenn Sie Fragen oder Kommentare haben, schreiben Sie mir an oder eröffnen Sie ein neues Thema in.
Dieser Code wurde inspiriert von den wunderbaren Artikeln:
In Zukunft werde ich weiterhin über Implementierungen von Hash-Tabellen für Grafikkarten schreiben und deren Leistung analysieren. Ich plane, verkettetes Hashing, Robin-Hood-Hashing und Cuckoo-Hashing unter Verwendung atomarer Operationen in Datenstrukturen zu behandeln, die für Grafikkarten geeignet sind.
Quelle: habr.com
