Welche Probleme bieten sich für GPU-Computing an?

84

Ich habe also ein gutes Gespür dafür, mit welchen Problemen ich am besten in Serie arbeite und welche parallel verwaltet werden können. Aber im Moment habe ich keine Ahnung, was bei CPU-basierten Berechnungen am besten funktioniert und was auf eine GPU ausgelagert werden sollte.

Ich weiß, dass es eine grundlegende Frage ist, aber ein Großteil meiner Suche wird von Leuten abgefangen, die sich eindeutig für das eine oder andere einsetzen, ohne wirklich zu rechtfertigen, warum oder etwas vage Faustregeln. Suchen Sie hier nach einer nützlicheren Antwort.

Fomite
quelle

Antworten:

63

GPU-Hardware hat zwei besondere Stärken: Raw Computing (FLOPs) und Speicherbandbreite. Die schwierigsten Rechenprobleme fallen in eine dieser beiden Kategorien. Beispielsweise liegt die Dichte der linearen Algebra (A * B = C oder Lösen [Ax = y] oder Diagonalisieren [A] usw.) abhängig von der Systemgröße irgendwo im Spektrum der Berechnungs- / Speicherbandbreite. Schnelle Fourier-Transformationen (FFT) erfüllen auch diese Anforderungen an die Gesamtbandbreite. Wie andere Transformationen, Gitter- / Gitter-basierte Algorithmen, Monte Carlo usw. Wenn Sie sich die NVIDIA SDK- Codebeispiele ansehen , können Sie ein Gefühl für die Art von Problemen bekommen, die am häufigsten angesprochen werden.

Ich denke, die lehrreichere Antwort ist die Frage: "In welchen Problemen sind GPUs wirklich schlecht?" Die meisten Probleme, die nicht in diese Kategorie fallen, können auf der GPU ausgeführt werden, obwohl einige mehr Aufwand erfordern als andere.

Probleme, die sich nicht gut abbilden lassen, sind im Allgemeinen zu klein oder zu unvorhersehbar. Bei sehr kleinen Problemen fehlt die Parallelität, die erforderlich ist, um alle Threads auf der GPU zu verwenden, und / oder sie können in einen Cache auf niedriger Ebene auf der CPU passen, wodurch die CPU-Leistung erheblich gesteigert wird. Unvorhersehbare Probleme weisen zu viele wichtige Verzweigungen auf, die ein effizientes Streaming von Daten vom GPU-Speicher zu den Kernen verhindern oder die Parallelität verringern können, indem das SIMD- Paradigma gebrochen wird (siehe „ Abweichende Verzerrungen “). Beispiele für solche Probleme sind:

  • Die meisten Grafikalgorithmen (zu unvorhersehbar, besonders im Speicherbereich)
  • Sparsame lineare Algebra (aber das ist auch schlecht für die CPU)
  • Kleine Signalverarbeitungsprobleme (z. B. FFTs mit weniger als 1000 Punkten)
  • Suche
  • Sortieren
Max Hutchinson
quelle
3
Dennoch GPU - Lösungen für die „unberechenbar“ Probleme sind möglich , und, während heute die Regel nicht möglich, in Zukunft an Bedeutung gewinnen kann.
linksum den
6
Ich möchte der Liste der GPU-Leistungsschalter speziell Zweige hinzufügen. Sie möchten, dass alle Ihre (Hunderte) denselben Befehl ausführen (wie in SIMD), um eine wirklich parallele Berechnung durchzuführen. Wenn beispielsweise auf AMD-Karten einer der Befehlsflüsse auf eine Verzweigung stößt und divergieren muss, divergieren alle Wellenfronten (parallele Gruppen). Wenn andere Einheiten von der Wellenfront nicht divergieren dürfen, müssen sie einen zweiten Durchgang durchführen. Das ist es, was maxhutch mit Vorhersehbarkeit meint, denke ich.
Violet Giraffe
2
@VioletGiraffe, das muss nicht unbedingt stimmen. In CUDA (dh auf Nvidia-GPUs) wirkt sich die Verzweigungsdivergenz nur auf die aktuelle Verzerrung aus, die höchstens 32 Threads beträgt. Obwohl verschiedene Warps denselben Code ausführen, sind sie nur dann synchron, wenn sie explizit synchronisiert werden (z __synchtreads(). B. mit ).
Pedro
1
@Pedro: Stimmt, aber das Verzweigen im Allgemeinen beeinträchtigt die Leistung. Bei Hochleistungscodes (welcher GPU-Code ist das nicht?) Ist dies fast unabdingbar.
Jvriesem
21

Probleme, die eine hohe arithmetische Intensität und regelmäßige Speicherzugriffsmuster aufweisen, sind in der Regel auf GPUs leicht (ier) zu implementieren und funktionieren auf diesen gut.

Die grundlegende Schwierigkeit bei der Verwendung von Hochleistungs-GPU-Code besteht darin, dass Sie über eine Tonne Kerne verfügen und möchten, dass alle Kerne so weit wie möglich ausgenutzt werden. Probleme mit unregelmäßigen Speicherzugriffsmustern oder nicht hoher Rechenintensität erschweren dies: Entweder verbringen Sie viel Zeit damit, Ergebnisse zu kommunizieren, oder Sie verbringen viel Zeit damit, Daten aus dem Speicher abzurufen (was langsam ist!), Und nicht genug Zeit, um Zahlen zu berechnen. Natürlich ist das Potenzial für Parallelität in Ihrem Code von entscheidender Bedeutung, damit er auch auf einer GPU gut implementiert werden kann.

Reid. Atcheson
quelle
Können Sie angeben, was Sie mit regulären Speicherzugriffsmustern meinen?
Fomite
1
Die Antwort von Maxhutch ist besser als meine. Was ich unter regulärem Zugriffsmuster verstehe, ist, dass auf den Speicher zeitlich und räumlich lokal zugegriffen wird. Das heißt: Sie machen nicht wiederholt große Gedächtnissprünge. Es ist auch so etwas wie ein Pauschalangebot, das mir aufgefallen ist. Es wird auch so verstanden, dass Ihre Datenzugriffsmuster entweder vom Compiler oder von Ihnen, dem Programmierer, vorgegeben werden können, so dass Verzweigungen (bedingte Anweisungen im Code) minimiert werden.
Reid.Atcheson
15

Dies ist nicht als eigenständige Antwort gedacht, sondern als Ergänzung zu den anderen Antworten von maxhutch und Reid.Atcheson .

Um das Beste aus GPUs herauszuholen, muss Ihr Problem nicht nur hoch (oder massiv) parallel sein, sondern auch der Kernalgorithmus, der auf der GPU ausgeführt wird, sollte so klein wie möglich sein. In OpenCL- Begriffen wird dies meist als Kernel bezeichnet .

Genauer gesagt sollte der Kernel in das Register jeder Multiprozessor-Einheit (oder Recheneinheit ) der GPU passen . Die genaue Größe des Registers ist abhängig von der GPU.

Da der Kernel klein genug ist, müssen die Rohdaten des Problems in den lokalen Speicher der GPU passen (gelesen: lokaler Speicher (OpenCL) oder gemeinsamer Speicher (CUDA) einer Recheneinheit). Andernfalls ist selbst die hohe Speicherbandbreite der GPU nicht schnell genug, um die Verarbeitungselemente ständig zu beschäftigen.
Normalerweise ist dieser Speicher etwa 16 bis 32 KiByte groß .

Torbjörn
quelle
Wird der lokale / gemeinsam genutzte Speicher jeder Prozessoreinheit nicht von allen Dutzenden (?) Threads gemeinsam genutzt, die in einem einzelnen Cluster von Kernen ausgeführt werden? Müssen Sie in diesem Fall den Arbeitsdatensatz nicht erheblich verkleinern, um die volle Leistung der GPU zu erzielen?
Dan Neely
Der lokale / gemeinsam genutzte Speicher einer Verarbeitungseinheit ist nur für die Recheneinheit selbst zugänglich und wird daher nur von den Verarbeitungselementen dieser Recheneinheit gemeinsam genutzt. Auf den globalen Speicher der Grafikkarte (in der Regel 1 GB) können alle Prozessoreinheiten zugreifen. Die Bandbreite zwischen den Verarbeitungselementen und dem lokalen / gemeinsam genutzten Speicher ist sehr schnell (> 1 TB / s), aber die Bandbreite zum globalen Speicher ist viel langsamer (~ 100 GB / s) und muss von allen Recheneinheiten gemeinsam genutzt werden.
Torbjörn,
Ich habe nicht nach dem Hauptspeicher der GPU gefragt. Ich dachte, der On-Die-Speicher wurde nur auf der Ebene des Kernclusters zugewiesen, nicht für jeden einzelnen Kern. ex für eine nVidia GF100 / 110 gpu; Für jeden der 16 SM-Cluster nicht die 512 Cuda-Kerne. Bei jedem SM, der für die parallele Ausführung von bis zu 32 Threads ausgelegt ist, muss zur Maximierung der GPU-Leistung der Arbeitssatz im Bereich von 1 KBit / Thread gehalten werden.
Dan Neely
@Torbjoern Sie möchten, dass alle GPU-Ausführungs-Pipelines ausgelastet sind. GPUs erreichen dies auf zwei Arten: (1) Die häufigste Methode besteht darin, die Belegung zu erhöhen oder anders auszudrücken, indem Sie die Anzahl der gleichzeitigen Threads erhöhen (kleine Kernel verwenden weniger von Threads) die freigegebenen Ressourcen, damit Sie mehr aktive Threads haben können); Vielleicht ist es besser, (2) die Parallelität der Befehlsebenen in Ihrem Kernel zu erhöhen, damit Sie einen größeren Kernel mit relativ geringer Belegung haben können (geringe Anzahl aktiver Threads). Siehe bit.ly/Q3KdI0
fcruz
11

Wahrscheinlich eine technischere Ergänzung zu den vorherigen Antworten: CUDA-GPUs (dh Nvidia) können als eine Reihe von Prozessoren beschrieben werden, die autonom mit jeweils 32 Threads arbeiten. Die Threads in jedem Prozessor arbeiten im Lock-Step (denken Sie an SIMD mit Vektoren der Länge 32).

Obwohl die verlockendste Art, mit GPUs zu arbeiten, darin besteht, vorzutäuschen, dass absolut alles im Gleichschritt abläuft, ist dies nicht immer die effizienteste Art, Dinge zu tun.

Wenn Ihr Code nicht ordnungsgemäß / automatisch mit Hunderten / Tausenden von Threads parallelisiert wird, können Sie ihn möglicherweise in einzelne asynchrone Tasks aufteilen, die sich gut parallelisieren lassen, und diejenigen mit nur 32 Threads ausführen, die im Sperrschritt ausgeführt werden. CUDA bietet eine Reihe atomarer Anweisungen, mit denen Mutexe implementiert werden können, die es den Prozessoren wiederum ermöglichen, sich untereinander zu synchronisieren und eine Liste von Aufgaben in einem Thread-Pool- Paradigma zu verarbeiten. Ihr Code würde dann genauso funktionieren wie auf einem Multi-Core-System. Denken Sie jedoch daran, dass jeder Core dann 32 Threads für sich hat.

Hier ist ein kleines Beispiel mit CUDA, wie dies funktioniert

/* Global index of the next available task, assume this has been set to
   zero before spawning the kernel. */
__device__ int next_task;

/* We will use this value as our mutex variable. Assume it has been set to
   zero before spawning the kernel. */
__device__ int tasks_mutex;

/* Mutex routines using atomic compare-and-set. */
__device__ inline void cuda_mutex_lock ( int *m ) {
    while ( atomicCAS( m , 0 , 1 ) != 0 );
    }
__device__ inline void cuda_mutex_unlock ( int *m ) {
    atomicExch( m , 0 );
    }

__device__ void task_do ( struct task *t ) {

    /* Do whatever needs to be done for the task t using the 32 threads of
       a single warp. */
    }

__global__ void main ( struct task *tasks , int nr_tasks ) {

    __shared__ task_id;

    /* Main task loop... */
    while ( next_task < nr_tasks ) {

        /* The first thread in this block is responsible for picking-up a task. */
        if ( threadIdx.x == 0 ) {

            /* Get a hold of the task mutex. */
            cuda_mutex_lock( &tasks_mutex );

            /* Store the next task in the shared task_id variable so that all
               threads in this warp can see it. */
            task_id = next_task;

            /* Increase the task counter. */
            next_tast += 1;

            /* Make sure those last two writes to local and global memory can
               be seen by everybody. */
            __threadfence();

            /* Unlock the task mutex. */
            cuda_mutex_unlock( &tasks_mutex );

            }

        /* As of here, all threads in this warp are back in sync, so if we
           got a valid task, perform it. */
        if ( task_id < nr_tasks )
            task_do( &tasks[ task_id ] );

        } /* main loop. */

    }

Sie müssen dann den Kernel mit aufrufen main<<<N,32>>>(tasks,nr_tasks), um sicherzustellen, dass jeder Block nur 32 Threads enthält und somit in eine einzelne Verzerrung passt. In diesem Beispiel habe ich der Einfachheit halber auch angenommen, dass die Aufgaben keine Abhängigkeiten haben (z. B. hängt eine Aufgabe von den Ergebnissen einer anderen ab) oder Konflikte (z. B. Arbeiten mit demselben globalen Speicher). Wenn dies der Fall ist, wird die Aufgabenauswahl etwas komplizierter, aber der Aufbau ist im Wesentlichen der gleiche.

Dies ist natürlich komplizierter als nur die Durchführung einer großen Anzahl von Zellen, erweitert jedoch die Art der Probleme, für die GPUs verwendet werden können, erheblich.

Pedro
quelle
2
Dies ist technisch gesehen richtig, es ist jedoch eine hohe Parallelität erforderlich, um eine hohe Speicherbandbreite zu erhalten, und die Anzahl der asynchronen Kernel-Aufrufe ist begrenzt (derzeit 16). Sie ist auch Tonnen von undokumentiertem Verhalten im Zusammenhang mit der Planung in der aktuellen Version. Ich würde raten, sich vorerst nicht auf asynchrone Kernel zu verlassen, um die Leistung zu verbessern ...
Max Hutchinson
2
Was ich beschreibe, kann alles in einem einzigen Kernelaufruf erledigt werden. Sie können N Blöcke mit jeweils 32 Threads erstellen, sodass jeder Block in einen Warp passt. Jeder Block erhält dann eine Aufgabe aus einer globalen Aufgabenliste (Zugriffskontrolle über Atomics / Mutexes) und berechnet sie mit 32 schrittweisen Threads. All dies geschieht in einem einzigen Kernel-Aufruf. Wenn Sie ein Codebeispiel möchten, lassen Sie es mich wissen, und ich werde eines veröffentlichen.
Pedro
4

Ein Punkt, der bisher nicht erwähnt wurde, ist, dass die aktuelle Generation von GPUs bei Gleitkommaberechnungen mit doppelter Genauigkeit nicht so gut abschneidet wie bei Berechnungen mit einfacher Genauigkeit. Wenn Ihre Berechnungen mit doppelter Genauigkeit ausgeführt werden müssen, können Sie davon ausgehen, dass sich die Laufzeit gegenüber der einfachen Genauigkeit um den Faktor 10 erhöht.

Brian Borchers
quelle
Ich möchte nicht zustimmen. Die meisten (oder alle) neueren GPUs bieten native Unterstützung für doppelte Genauigkeit. Fast jede solche GPU meldet Berechnungen mit doppelter Genauigkeit, die mit ungefähr der Hälfte der Geschwindigkeit der einfachen Genauigkeit ausgeführt werden, was wahrscheinlich auf die einfache Verdoppelung der erforderlichen Speicherzugriffe / Bandbreite zurückzuführen ist.
Godric Seer
1
Zwar bieten die neuesten und besten Nvidia Tesla-Karten eine Spitzenleistung mit doppelter Präzision, die nur halb so hoch ist wie die Spitzenleistung mit einfacher Präzision, aber das Verhältnis beträgt 8 zu 1 für häufigere Kundenkarten mit Fermi-Architektur.
Brian Borchers
@GodricSeer Das 2: 1-Verhältnis von SP und DP-Gleitkomma hat sehr wenig mit Bandbreite zu tun und fast alles hängt davon ab, wie viele Hardwareeinheiten zur Ausführung dieser Operationen vorhanden sind. Es ist üblich, die Registerdatei für SP und DP wiederzuverwenden, daher kann die Gleitkommaeinheit die SP-Operationen zweimal als DP-Operationen ausführen. Es gibt zahlreiche Ausnahmen zu diesem Design, z. B. IBM Blue Gene / Q (hat keine SP-Logik und daher läuft SP mit ~ 1.05x DP). Einige GPUs haben andere Verhältnisse als 2, z. B. 3 und 5.
Jeff
Es ist vier Jahre her, seit ich diese Antwort geschrieben habe, und die aktuelle Situation mit NVIDIA-GPUs ist, dass für die GeForce- und Quadro-Leitungen das DP / SP-Verhältnis jetzt 1/32 beträgt. NVIDIAs Tesla-GPUs bieten eine viel stärkere Double-Precision-Leistung, kosten aber auch viel mehr. Auf der anderen Seite hat AMD die Leistung seiner Radeon-GPUs bei doppelter Präzision nicht auf die gleiche Weise beeinträchtigt.
Brian Borchers
4

Aus metaphorischer Sicht kann die GPU als eine Person gesehen werden, die auf einem Nagelbett liegt. Die Person, die oben liegt, enthält die Daten und in der Basis jedes Nagels befindet sich ein Prozessor. Der Nagel ist also ein Pfeil, der vom Prozessor zum Speicher zeigt. Alle Nägel haben ein regelmäßiges Muster wie ein Gitter. Wenn der Körper gut verteilt ist, fühlt er sich gut an (Leistung ist gut), wenn der Körper nur einige Stellen des Nagelbettes berührt, sind die Schmerzen schlecht (schlechte Leistung).

Dies kann als ergänzende Antwort zu den hervorragenden Antworten oben verstanden werden.

labotsirc
quelle
4

Alte Frage, aber ich denke, dass diese Antwort von 2014 - bezogen auf statistische Methoden, aber verallgemeinerbar für jeden, der weiß, was eine Schleife ist - besonders anschaulich und informativ ist.

GT.
quelle
2

GPUs haben lange Latenzzeiten für E / A, daher müssen viele Threads verwendet werden, um den Speicher zu überlasten. Um einen Warp beschäftigt zu halten, sind viele Threads erforderlich. Wenn der Codepfad 10 Takte und die E / A-Latenz 320 Takte beträgt, sollten 32 Threads der Sättigung des Warps nahe kommen. Wenn der Codepfad 5 Takte beträgt, verdoppeln Sie die Threads.

Suchen Sie mit tausend Kernen nach Tausenden von Threads, um die GPU voll auszunutzen.

Der Speicherzugriff erfolgt über eine Cache-Zeile, normalerweise 32 Byte. Das Laden eines Bytes ist mit 32 Bytes vergleichbar. Verschmelzen Sie den Speicher, um die Nutzungslokalität zu erhöhen.

Es gibt viele Register und einen lokalen RAM für jeden Warp, so dass die Nachbarn sich teilen können.

Näherungssimulationen großer Mengen sollten sich gut optimieren lassen.

Zufällige I / O und Single Threading ist ein Vergnügen ...

user14381
quelle
Dies ist eine wirklich faszinierende Frage; Ich diskutiere mit mir selbst darüber, ob es möglich (oder die Mühe wert) ist, eine einigermaßen einfache Aufgabe (Kantenerkennung in Luftbildern) zu "parallelisieren", wenn jede Aufgabe ~ 0,06 Sekunden dauert, aber ~ 1,8 Millionen Aufgaben zu erledigen sind ( pro Jahr für Daten im Wert von 6 Jahren: Die Aufgaben sind definitiv trennbar) ... also ~ 7,5 Tage Rechenzeit auf einem Kern. Wenn jeder Berechnungsvorgang auf einer GPU schneller war und der Auftrag 1-pro-nGPU-Punkte [n klein] parallelisiert werden konnte, ist es tatsächlich wahrscheinlich, dass die Auftragszeit auf ~ 1 Stunde sinken könnte? Scheint unwahrscheinlich.
GT.
0

Stellen Sie sich ein Problem vor, das mit viel roher Gewalt gelöst werden kann, wie z. B. Travelling Salesman. Stellen Sie sich vor, Sie haben Server-Racks mit jeweils 8 verspielten Grafikkarten und jede Karte hat 3000 CUDA-Kerne.

Lösen Sie einfach ALLE möglichen Routen des Verkäufers und sortieren Sie sie nach Zeit / Entfernung / Metrik. Sicher, Sie werfen fast 100% Ihrer Arbeit weg, aber manchmal ist rohe Gewalt eine praktikable Lösung.

Criggie
quelle
Ich hatte eine Woche lang Zugriff auf eine kleine Farm mit 4 solchen Servern, und in fünf Tagen habe ich mehr Distributed.net-Blöcke erstellt als in den letzten 10 Jahren.
Criggie
-1

Wenn ich viele technische Ideen studiere, würde ich sagen, dass eine GPU eine Form der Fokussierung von Aufgaben, des Speichermanagements, der wiederholbaren Berechnung ist.

Viele Formeln sind zwar einfach zu schreiben, aber schwierig zu berechnen. In der Matrixmathematik erhalten Sie nicht eine einzige Antwort, sondern viele Werte.

Dies ist wichtig, wenn Sie berechnen, wie schnell ein Computer Werte berechnet und Formeln ausführt, da einige Formeln nicht ohne alle berechneten Werte ausgeführt werden können (daher langsamer). Ein Computer weiß nicht genau, in welcher Reihenfolge Formeln ausgeführt oder Werte berechnet werden müssen, um sie in diesen Programmen zu verwenden. Es zwingt sich hauptsächlich mit hoher Geschwindigkeit durch und zerlegt Formeln in Spannfutter, um zu berechnen, aber viele Programme erfordern heutzutage diese berechneten Spannfutter und warten in Ques (und Ques of Ques und mehr Ques of Ques).

Zum Beispiel in einem Simulationsspiel, das zuerst bei Kollisionen den Schaden der Kollision, die Position der Objekte, die neue Geschwindigkeit berechnet werden soll? Wie viel Zeit sollte das dauern? Wie kann eine CPU mit dieser Last umgehen? Außerdem sind die meisten Programme sehr abstrakt und erfordern mehr Zeit für die Datenverarbeitung. Sie sind nicht immer für Multithreading ausgelegt oder bieten in abstrakten Programmen keine guten Möglichkeiten, dies effektiv zu tun.

Als die CPU immer besser wurde, wurden die Leute beim Programmieren schlampig und wir mussten auch für viele verschiedene Computertypen programmieren. Eine GPU ist so konzipiert, dass sie durch viele einfache Berechnungen gleichzeitig brachiale Gewalt ausübt (Speicher (sekundär / RAM) und Heizung, Kühlung sind die wichtigsten Engpässe beim Rechnen). Eine CPU verwaltet viele viele Fragen gleichzeitig oder wird in viele Richtungen gezogen, um herauszufinden, was zu tun ist, wenn man nicht in der Lage ist, dies zu tun. (Hey, es ist fast menschlich)

Eine GPU ist Grunzarbeiter die mühsame Arbeit. Eine CPU verwaltet das gesamte Chaos und kann nicht mit jedem Detail umgehen.

Also, was lernen wir? Eine GPU erledigt mühsame Detailarbeiten auf einmal und eine CPU ist eine Multitask-Maschine, die sich nicht sehr gut auf zu viele Aufgaben konzentrieren kann. (Es ist wie es gleichzeitig Aufmerksamkeitsstörung und Autismus hat).

Engineering gibt es die Ideen, Design, Realität und viel Grunzarbeit.

Wenn ich gehe, denke daran, einfach zu beginnen, schnell zu beginnen, schnell zu scheitern, schnell zu scheitern und nie aufzuhören, es zu versuchen.

Andrew G. Corbi
quelle