Ich habe auf Github .
Das ist eine einfache Hash-Tabelle für GPUs, die in der Lage ist, hunderte Millionen Einfügungen pro Sekunde zu verarbeiten. Auf meinem Laptop mit NVIDIA GTX 1060 fügt der Code ungefähr 64 Millionen zufällig generierter Schlüssel-Wert-Paare in etwa 210 ms ein und entfernt 32 Millionen Paare in etwa 64 ms.
Das bedeutet, dass die Geschwindigkeit auf dem Laptop etwa 300 Millionen Einfügungen/Sekunde und 500 Millionen Löschungen/Sekunde beträgt.
Die Tabelle ist in CUDA geschrieben, obwohl dieselbe Methode auch auf HLSL oder GLSL angewendet werden kann. Die Implementierung hat mehrere Einschränkungen, die eine hohe Leistung auf der Grafikkarte gewährleisten:
- Es werden nur 32-Bit-Schlüssel und -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 , das heißt, es handelt sich einfach um ein Array von Schlüssel-Wert-Paaren, das im Speicher gespeichert ist und hervorragende Cache-Leistung bietet. Dies kann nicht von einer Kettenspeicherung (chaining) gesagt werden, die ein Suchen eines Zeigers in einer verketteten Liste impliziert. Eine 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 keine Primzahl, da 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 im Fall der linearen Erkundung, da beim Durchsuchen der Tabelle der Slot-Index in jeden Slot gewickelt werden muss. Folglich entsteht bei jedem Slot ein Kostenfaktor durch die Modulo-Operation.
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 angegebenen Code wird der Murmur3-Hash verwendet, der nur einige Verschiebungen, XOR-Operationen und Multiplikationen durchführt.
In der Hash-Tabelle wird eine Schutzmethodik gegen Blockaden angewendet, die unabhängig von der Reihenfolge in der Speichermmlung funktioniert. Selbst wenn einige Schreiboperationen die Reihenfolge anderer solcher Operationen stören, bleibt der Zustand der Hash-Tabelle dennoch korrekt. Darüber sprechen wir weiter unten. Diese Methode funktioniert hervorragend mit Grafikkarten, auf denen tausende von Threads gleichzeitig ausgeführt werden.
Die Schlüssel und Werte in der Hash-Tabelle werden leer initialisiert.
Der Code kann so modifiziert werden, dass er sowohl 64-Bit-Schlüssel als auch -Werte verarbeiten kann. Für Schlüssel sind atomare Lese-, Schreib- und Vergleichsoperationen mit Austausch (compare-and-swap) erforderlich. Für Werte sind atomare Lese- und Schreiboperationen notwendig. Glücklicherweise sind in CUDA 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 der vier Zustände haben:
- Der Schlüssel und der Wert sind leer. In diesem Zustand wird die Hashtabelle initialisiert.
- Der Schlüssel wurde geschrieben, der Wert jedoch noch nicht. Wenn zu diesem Zeitpunkt ein anderer Ausführungs-Thread die Daten liest, wird er anschließend einen leeren Wert zurückgeben. Das ist normal; das Gleiche würde passieren, wenn ein anderer Ausführungs-Thread etwas früher gearbeitet hätte, und wir sprechen hier von einer konkurrierenden Datenstruktur.
- Sowohl der Schlüssel als auch der Wert wurden geschrieben.
- Der Wert ist für andere Ausführungs-Threads verfügbar, der Schlüssel jedoch noch nicht. Das kann passieren, da das Programmiermodell in CUDA ein schwach geordnetes Modellspeichermodell vorsieht. Das ist in Ordnung; jedenfalls ist der Schlüssel weiterhin leer, auch wenn der Wert das nicht mehr 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, werden wir später darüber sprechen.
Der Code der Hashtabelle funktioniert auch mit schwach geordneten Speichermodellen, bei denen die Reihenfolge von Lese- und Schreibzugriffen auf den Speicher unbekannt ist. Wenn wir das Einfügen, Suchen und Löschen in der Hashtabelle besprechen, denken Sie daran, dass jedes Schlüssel-Wert-Paar sich in einem der vier oben beschriebenen Zustände befindet.
Einfügen in die Hash-Tabelle
Die CUDA-Funktion, die Paare aus Schlüssel und Wert in die 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, iteriert der Code über das Array der Hash-Tabelle, beginnend mit dem Hash des einzufügenden Schlüssels. In jedem Slot des Arrays wird eine atomare Vergleichsoperation mit Austausch durchgeführt, bei der der Schlüssel in diesem Slot mit leer verglichen wird. Wenn eine Nichtübereinstimmung festgestellt wird, wird der Schlüssel im Slot auf den einzufügenden Schlüssel aktualisiert, und der ursprüngliche Schlüssel des Slots wird zurückgegeben. Wenn dieser ursprüngliche Schlüssel leer war oder dem einzufügenden Schlüssel entsprach, hat der Code einen geeigneten Slot zum Einfügen gefunden und fügt den Wert in den Slot ein.
Wenn in einem Kernaufruf gpu_hashtable_insert() Wenn mehrere Elemente denselben Schlüssel haben, kann eines ihrer Werte in den Schlüssel-Slot geschrieben werden. Das gilt als normal: Eine der Schreiboperationen für den Schlüssel-Wert-Paar während des Aufrufs wird erfolgreich sein, aber da das alles parallel in mehreren Ausführungsthreads geschieht, können wir nicht vorhersagen, welche Schreiboperation im Speicher zuletzt erfolgt.
Suche in der Hash-Tabelle
Code zur 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 Schlüssels, der in der Tabelle gespeichert ist, zu finden, iterieren wir durch das Array, beginnend mit dem Hash des gesuchten Schlüssels. In jedem Slot überprü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 sofern dies der Fall ist, brechen wir die Suche ab.
Wenn wir den Schlüssel nicht finden können, gibt der Code einen leeren Wert zurück.
Alle diese Suchoperationen können während Einfügungen und Löschungen gleichzeitig durchgeführt werden. Jedes Paar in der Tabelle hat für den Thread eines 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 dem Unterschied, dass beim Erkennen einer Übereinstimmung mit dem Schlüssel sein Wert leer gemacht wird.
Wie bereits erwähnt, sobald der Schlüssel in einen Slot geschrieben wurde, wird er nicht mehr bewegt. 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 Schreiboperation für den Wert des Slots verwenden müssen, denn es spielt keine Rolle, ob der aktuelle Wert leer ist oder nicht – er wird sowieso leer werden.
Vergrößern der Hash-Tabelle
Die Größe der Hashtabelle kann durch die Erstellung einer größeren Tabelle und das Einfügen von nicht-leeren Elementen aus der alten Tabelle geändert werden. Diese Funktionalität habe ich nicht implementiert, da ich den Code einfach halten wollte. Darüber hinaus erfolgt in CUDA-Programmen die Speicherzuweisung häufig im Host-Code und nicht im CUDA-Kernel.
In dem Artikel beschreibt, wie man eine solche gegen Blockierungen geschützte Datenstruktur ändert.
Konkurrierende Prozesse
In den oben genannten Codefragmenten gpu_hashtable_insert(), _lookup() und _delete() verarbeiten jeweils ein Schlüssel-Wert-Paar zurzeit. Unten gpu_hashtable_insert(), _lookup() und _delete() verarbeiten sie ein Array von Paaren parallel, wobei jedes Paar in einem separaten GPU-Thread verarbeitet wird:
// 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 blockierungsfreie Hashtabelle unterstützt konkurrierende Einfügungen, Suchen und Löschungen. Da sich Schlüssel-Wert-Paare immer in einem der vier Zustände befinden und die Schlüssel sich nicht bewegen, gewährleistet die Tabelle Korrektheit selbst bei gleichzeitiger Durchführung unterschiedlicher Operationen.
Wenn wir jedoch gleichzeitig Verarbeitungspakete für Einfügungen und Löschungen haben und im Eingabearray Schlüsselpaare dupliziert sind, können wir nicht vorhersagen, welche Paare "gewinnen" werden — die zuletzt in die Hash-Tabelle geschriebenen. Angenommen, wir haben den Einfügecode mit dem Eingabearray 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 sicher in der Tabelle vorhanden sein, aber auch eines der Paare A/0, A/2 oder A/4. Das kann ein Problem sein, kann aber auch keins sein – es hängt alles von der Anwendung ab. Möglicherweise wissen Sie im Voraus, dass im Eingabearray keine doppelten Schlüssel vorhanden sind, oder es ist Ihnen egal, welches Wert zuletzt gespeichert wurde.
Wenn dies für Sie ein Problem darstellt, müssen Sie die doppelten Paare auf verschiedene systemseitige 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/3aufgerufen wird, 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() Verwenden Sie einen einfachen (plain) oder variablen (volatile) Zeiger auf ein Array von Paaren in einer Hash-Tabelle. stellt fest, dass:
Der Compiler kann nach eigenem Ermessen Lese- und Schreiboperationen im globalen oder gemeinsamen Speicher optimieren … Diese Optimierungen können mit dem Schlüsselwort
volatile: … jeder Verweis auf diese Variable wird zu einer tatsächlichen Lese- oder Schreibanweisung im Speicher kompiliert.
Überlegungen zur Korrektheit erfordern keine Anwendung volatile. Wenn der Ausführungsfaden einen zwischengespeicherten Wert aus einer früheren Leseoperation verwendet, bedeutet dies, dass er leicht veraltete Informationen verwendet. Dennoch handelt es sich um Informationen aus dem korrekten 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: Laut meinen Tests verringerte sich bei der Löschung von 32 Millionen Elementen die Geschwindigkeit von 500 Millionen Löschungen/Sekunde auf 450 Millionen Löschungen/Sekunde.
Leistung
Im Test zum Einfügen von 64 Millionen Elementen und dem Entfernen von 32 Millionen davon besteht praktisch keine Konkurrenz zwischen std::unordered_map und der Hash-Tabelle für die GPU.
std::unordered_map hat 70.691 ms für das Einfügen und Löschen von Elementen mit anschließender Freigabe benötigt unordered_map (die Freigabe von Millionen von Elementen dauert eine geraume Zeit, da innerhalb unordered_map zahlreiche Speicherzuweisungen durchgeführt werden). Ehrlich gesagt haben std:unordered_map ganz andere Einschränkungen. Es handelt sich um einen einzelnen CPU-Thread, der Schlüssel-Werte beliebiger Größe unterstützt, gut bei hohen Ausnutzungsgraden funktioniert und nach häufigen Löschvorgängen eine stabile Leistung zeigt.
Die Laufzeit der Hash-Tabelle für GPU und Interprozesskommunikation betrug 984 ms. Dies umfasst die Zeit, die für die Platzierung der Tabelle im Speicher und ihre Löschung (eine einmalige Zuweisung von 1 GB Speicher, die in CUDA einige Zeit in Anspruch nimmt), das Einfügen und Löschen von Elementen sowie deren Iteration benötigt wird. Auch alle Kopiervorgänge in und aus dem Speicher der Grafikkarte sind berücksichtigt.
Die Funktion der Hash-Tabelle benötigte 271 ms. Dazu zählt die Zeit, die die Grafikkarte für das Einfügen und Löschen von Elementen benötigt, jedoch nicht die Zeit für das Kopieren in den Speicher und das Durchlaufen der resultierenden Tabelle. Wenn die GPU-Tabelle lange lebt oder die Hash-Tabelle vollständig im Grafikspeicher enthalten ist (zum Beispiel für die Erstellung einer Hash-Tabelle, die von anderem GPU-Code und nicht vom Prozessor verwendet wird), ist das Testergebnis relevant.
Die Hash-Tabelle für die Grafikkarte zeigt eine hohe Leistung dank der großen Bandbreite und aktiven Parallelisierung.
Nachteile
Die Architektur der Hash-Tabelle hat einige Probleme, die man beachten sollte:
- Die lineare Erfassung wird durch die Clusterbildung behindert, wodurch die Schlüssel in der Tabelle nicht optimal platziert sind.
- Schlüssel werden nicht mit der Funktion entfernt
löschenund verstopfen im Laufe 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 durchführt. Eine Möglichkeit, diese Nachteile zu mildern, besteht darin, in eine neue Tabelle mit einem ausreichenden niedrigen Auslastungsgrad neu zu hashen und entfernte Schlüssel während des Neuhashens zu filtern.
Um die beschriebenen Probleme zu veranschaulichen, nutze ich den oben genannten Code zur Erstellung einer Tabelle mit 128 Millionen Elementen und werde zyklisch 4 Millionen Elemente einfügen, bis 124 Millionen Slots belegt sind (Auslastungsgrad von ungefähr 0,96). Hier ist die Ergebnistabelle, jede Zeile ist ein Aufruf des CUDA-Kernels mit der Einfügung von 4 Millionen neuen Elementen in eine Hash-Tabelle:
Auslastungsgrad
Dauer der Einfügung von 4 194 304 Elementen
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 sinkt die Leistung, was in den meisten Fällen unerwünscht ist. Wenn eine Anwendung Elemente in eine Tabelle einfügt und diese dann wieder verwirft (z.B. beim Zählen von Wörtern in einem Buch), ist das kein Problem. Sollte jedoch eine Anwendung eine langlebige Hash-Tabelle verwenden (z.B. in einem Grafikeditor zum Speichern nicht leerer Bildteile, wenn der Benutzer häufig Informationen einfügt und löscht), kann dieses Verhalten problematisch sein.
Ich habe die Tiefe der Hash-Tabellen-Sondierung nach 64 Millionen Einfügungen (Auslastung 0,5) gemessen. Die durchschnittliche Tiefe betrug 0,4774, sodass sich die meisten Schlüssel entweder in dem besten verfügbaren Slot oder nur einen Slot von der besten Position entfernt befanden. Die maximale Sondierungstiefe betrug 60.
Anschließend maß ich die Sondierungstiefe in einer Tabelle mit 124 Millionen Einfügungen (Auslastung 0,97). Die durchschnittliche Tiefe betrug bereits 10,1757, während die maximale - 6474 (!!). Die Leistung der linearen Sondierung sankt erheblich bei hohen Auslastungsgraden.
Es ist am besten, die Auslastung dieser Hash-Tabelle niedrig zu halten. Aber das führt dann zu einer höheren Leistung auf Kosten des Speicherverbrauchs. Glücklicherweise kann dies bei 32-Bit-Schlüsseln und -Werten gerechtfertigt sein. Wenn wir im obigen Beispiel bei einer Tabelle mit 128 Millionen Elementen eine Auslastung von 0,25 beibehalten, können wir höchstens 32 Millionen Elemente unterbringen, was bedeutet, dass 96 Millionen Slots verloren gehen — jeweils 8 Byte pro Paar, also 768 MB verlorenem Speicher.
Beachten Sie, dass es hier um den Speicherverlust der Grafikkarte geht, der wertvoller 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 unklug, solche Mengen zu verlieren.
Später werde ich detaillierter über die Erstellung von Hash-Tabellen für Grafikkarten schreiben, die keine Probleme mit der Durchdringungstiefe haben, sowie über Möglichkeiten zur Wiederverwendung verlorener Slots.
Messung der Durchdringungstiefe
Um die Tiefe der Schlüsselabfrage zu bestimmen, können wir den Hash des Schlüssels (seinen idealen Index in der Tabelle) aus seinem tatsächlichen Tabelleneintrag extrahieren:
// 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ärer Zahlen im Zweierkomplement und der Tatsache, dass die Kapazität der Hashtabelle gleich der Potenz von zwei ist, wird dieser Ansatz sogar dann funktionieren, wenn der Schlüsselsatz an den Anfang der Tabelle verschoben wird. Nehmen wir einen Schlüssel, der in 1 gehasht wird, aber in Steckplatz 3 eingefügt wurde. Für eine Tabelle mit einer Kapazität von 4 erhalten wir dann (3 — 1) & 3, was 2 entspricht.
Fazit
Wenn Sie Fragen oder Kommentare haben, kontaktieren Sie mich unter oder eröffnen Sie ein neues Thema in.
Dieser Code ist inspiriert von den wunderbaren Artikeln:
In Zukunft werde ich weiterhin über Implementierungen von Hashtabellen für Grafikkarten schreiben und deren Leistung analysieren. Ich plane Kettenbindung, Robin-Hood-Hashing und Cuckoo-Hashing unter Verwendung atomarer Operationen in für Grafikkarten geeigneten Datenstrukturen.
Quelle: habr.com
