Sortieralgorithmus mit CUDA

(ashwanirathee.com)

3 Punkte von GN⁺ 2025-03-13 | 2 Kommentare | Auf WhatsApp teilen

So lässt sich mit CUDA über Parallel Computing die Performance eines Sortieralgorithmus verbessern
Eine grundlegende Merge-Sort-Implementierung wird in CUDA umgesetzt, um das Potenzial für Performance-Verbesserungen zu untersuchen

Grundlegender rekursiver Merge Sort (CPU-Implementierung)

Ein Sortieralgorithmus, der ein Array in zwei Teilarrays aufteilt, diese jeweils sortiert und anschließend zusammenführt
Das Array wird rekursiv aufgeteilt; sobald die Größe 1 erreicht, wird der Merge-Schritt ausgeführt
Wichtige Punkte zur Implementierung
- Verwendung von uint8_t → minimaler Speicherverbrauch für kleine Werte (0–255)
- Verwendung von long long → Verarbeitung großer Arrays (bis zu 10¹⁸) möglich
- Zur Validierung der Ergebnisse für den Performance-Vergleich wird std::sort verwendet
- Zeitkomplexität: durchschnittlich O(n log n)
- Speicherkomplexität: O(n)

Grundlegender rekursiver Merge Sort in CUDA

Folgt demselben Muster wie die CPU-Implementierung
Der Merge-Schritt ist so implementiert, dass er in CUDA parallel ausgeführt wird
Wichtige Punkte zur Implementierung
- Verwendung von cudaMalloc, cudaMemcpy, cudaFree → GPU-Speicherzuweisung und Datenübertragung
- merge<<<1, 1>>>(...) → führt den Merge-Schritt parallel aus
- cudaDeviceSynchronize() → Synchronisation bis zum Abschluss des Merges
- Performance-Problem → CUDA ist für rekursive Verarbeitung ineffizient, daher ist ein iterativer Ansatz nötig

Vergleich von CPU- und GPU-Implementierung

Da rekursive Aufrufe auf der CPU ausgeführt werden, kommt es zu Performance-Verlusten
Rekursive Aufrufe in CUDA verursachen Probleme mit der Stack-Größe und Kernel-Launch-Overhead
Methode zur Performance-Verbesserung: Umstellung auf einen iterativen Bottom-up-Ansatz

Iterativer Bottom-up-Merge Sort (CPU-Implementierung)

Beginn mit kleinen Teilarrays und schrittweises Zusammenführen → in CUDA effizienter
Die Größe der zu mergenden Arrays wird in den Schritten 1, 2, 4, 8, … erhöht

Wichtige Codestruktur

MERGE_SORT(arr, temp, start, end)  
  FOR sub_size ← 1 TO end STEP 2 × sub_size DO  
      FOR left ← 0 TO end STEP 2 × sub_size DO  
          mid ← MIN(left + sub_size - 1, end)  
          right ← MIN(left + 2 * sub_size - 1, end)  
          MERGE(arr, temp, left, mid, right)  
      ENDFOR  
  ENDFOR  
END MERGE_SORT

Wichtige Punkte zur Implementierung
- Falls die Array-Größe kein Vielfaches von 2 ist, wird das Problem durch Clamping der Indizes gelöst
- Der Merge-Schritt wird in Schleifen ausgeführt
- Großes Potenzial für Performance-Verbesserungen

Iterativer Bottom-up-Merge Sort (CUDA-Implementierung)

Performance-Verbesserung durch parallele Ausführung des iterativen Merge Sort
Zur parallelen Ausführung der Merge-Schritte werden Anzahl der Threads und Blöcke berechnet und dann gestartet

Wichtige Codestruktur

  void mergeSort(uint8_t* arr, uint8_t* temp, long long n) {  
      bool flipflop = true;  
      long long size;  
      for (size = 1; size < n; size *= 2) {  
          numThreads = max(n / (2 * size), (long long)1);  
          gridSize = (numThreads + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;  
          mergeKernel<<<gridSize, THREADS_PER_BLOCK>>>(flipflop ? arr : temp, flipflop ? temp : arr, size, n);  
          CUDA_CHECK(cudaGetLastError());  
          CUDA_CHECK(cudaDeviceSynchronize());  
          flipflop = !flipflop;  
      }  
      if (!flipflop) CUDA_CHECK(cudaMemcpy(arr, temp, n * sizeof(uint8_t), cudaMemcpyDeviceToDevice));  
  }

Wichtige Punkte
- flipflop → Umschalten des Speicherorts für das Merge-Ergebnis
- gridSize, THREADS_PER_BLOCK → Parallelisierung der Merge-Aufgaben
- mergeKernel → weist jedem Thread eine eigene Merge-Aufgabe zu
- Indexverwaltung durch Berechnung von Thread- und Block-Indizes

Performance-Ergebnisse

Bottom-up-Merge Sort (CUDA) → klare Performance-Verbesserung
- Kleine Arrays → CPU ist schneller
- Große Arrays → CUDA hat einen Performance-Vorteil
thrust::sort → starke GPU-Performance bei großen Arrays
Die Performance-Verbesserung von CUDA wird durch den Overhead der Datenübertragung begrenzt

Fazit und zukünftige Arbeiten

Erfolgreiche Performance-Verbesserung eines CUDA-basierten Merge Sort
Wichtige Lernergebnisse:
- Verständnis der CUDA-Konzepte für parallele Verarbeitung und Strategien zum Performance-Tuning
- Iterativer Merge Sort → in CUDA effektiver als ein rekursiver Ansatz
- CUDA-spezifische Performance-Engpässe wie Thread-Synchronisation und Speicherübertragung identifiziert
Künftige Verbesserungen:
- Aufgabentrennung und Optimierung zwischen CPU und GPU
- Performance-Tests mit noch größeren Arrays
- Kombination von thrust::sort mit benutzerdefiniertem Implementierungscode
- Performance-Optimierung durch Nutzung von Shared Memory

2 Kommentare

kandk 2025-03-14

Das ist wohl mit Radix Sort in CUDA implementiert; ich habe die Erfahrung gemacht, es anhand der Referenz genauso umzusetzen.

GN⁺ 2025-03-13

Hacker-News-Kommentare

Es geht nicht darum, wie man auf der GPU schnell sortiert. Der schnellste Algorithmus in CUDA ist Onesweep; er nutzt die Parallelität der GPU und verwendet komplexe Techniken, um Einschränkungen zu überwinden.
- Linebender arbeitet daran, diese Ideen portabler auf GPUs anzuwenden.
- Weiteres Material dazu findet sich auf der Wiki-Seite von Linebender.
Wie auch andere Kommentare anmerken, ist dieser Algorithmus nicht geeignet. Algorithmen wie Onesweep sind cool, aber schwer zu verstehen.
- Wenn man sich den zugrunde liegenden Algorithmus, also Radix Sort, ansieht, ist es leichter nachvollziehbar.
- Radix Sort lässt sich sehr einfach parallelisieren und ist ein schöner, eleganter Ansatz.
Als kleines Spielzeugproblem ist das interessant. Mit Optionen zur Thread-Abstimmung könnte man die Performance verbessern.
- Es macht auch Spaß, mit Nsight herauszufinden, welche Faktoren die Performance beeinflussen.
- Man sollte auch andere Sortieralgorithmen in Betracht ziehen. Netzwerk-Sortierverfahren wie Bitonic Sort erfordern zwar mehr Arbeit, können auf paralleler Hardware aber schneller laufen.
- Ich habe auf einer H100 eine einfache Implementierung gebaut, die 10M in etwa 10 ms sortiert. Mit mehr Arbeit könnte sie noch schneller werden.
Die Sprache Futhark macht solche Algorithmen auf der GPU bequemer nutzbar.
- Es ist eine sehr hochsprachliche Sprache, die zu GPU-Befehlen kompiliert und als Python-Bibliothek zugänglich ist.
- Auf der Website gibt es ein Beispiel für eine Merge-Sort-Implementierung.
Das erinnert mich an ein kleines Projekt aus meiner Studienzeit, in dem ich Bitonic Sort in CUDA implementiert habe.
- Die Bitonic-Sort-Implementierung ist auf GitHub zu finden.
Die Notizen, die das Konzept der GPU-Thread-Indizierung erklären, sind gut.
- Es wird ein persönlicher Blogpost zu den Performance-Vorteilen vektorisierter Sortierung vorgestellt.
Ich mag die schnelle Radix-Sort-Implementierung.
- Sie arbeitet problemlos mit der Cuda Driver API und ist im Gegensatz zu CUB nicht auf die Runtime API beschränkt.
- Onesweep ist ebenfalls enthalten, aber ich konnte es nicht zum Laufen bringen.
Ich wollte das zusammen mit Unity verwenden, konnte den Flaschenhals beim Datentransfer aber nicht überwinden.
- Auch bei der Verwendung von Compute Shadern gibt es Overhead, aber nicht besonders viel.
Damit sich Sortieren auf der GPU lohnt, braucht man große Arrays.
- Die Datenübertragung zwischen RAM und GPU kostet Zeit.
Um Zeit zu sparen, die Zusammenfassung: Jemand hat einen Sortieralgorithmus für die GPU geschrieben, aber er war langsam.
- Es ist nicht State of the Art, und es ist unklar, ob der Autor weiß, wie man die GPU effektiv nutzt.
- Es ist nur ein persönliches Experiment mit GPU-Programmierung.