Kapitel 11: GPU-Forschungsrichtungen zur Skalierung und affinen Ausführung
Wie in Kapitel 2 beschrieben, bieten GPU-Computing-APIs wie CUDA und OpenCL ein MIMD-ähnliches Programmiermodell, das es dem Programmierer ermöglicht, eine große Anzahl von skalaren Threads auf die GPU zu starten. Während jeder dieser skalaren Threads seinen eigenen Ausführungspfad verfolgen und beliebige Speicherorte zugreifen kann, folgen sie in den meisten Fällen einer kleinen Anzahl von Ausführungspfaden und führen ähnliche Operationen durch.
Die konvergente Kontrollfluss-Struktur unter GPU-Threads wird in den meisten, wenn nicht sogar allen, modernen GPUs über das SIMT-Ausführungsmodell ausgenutzt, bei dem skalare Threads in Warps gruppiert werden, die auf SIMD-Hardware ausgeführt werden (siehe Abschnitt 3.1.1). Dieses Kapitel fasst eine Reihe von Forschungsarbeiten zusammen, die die Ähnlichkeit dieser skalaren Threads durch Skalierung und affine Ausführung weiter ausnutzen.
Der Schlüsseleinblick dieser Forschung liegt in der Beobachtung der Wertstruktur [Kim et al., 2013] über Threads hinweg, die denselben Compute-Kernel ausführen. Die beiden Arten von Wertstruktur, uniform und affin, werden im Compute-Kernel in Beispiel 11.1 veranschaulicht.
Uniform-Variable
Eine Variable, die denselben konstanten Wert für jeden Thread im Compute-Kernel hat. In Algorithmus 11.1 haben die Variable a sowie die Literale THRESHOLD und Y_MAX_VALUE alle uniforme Werte über alle Threads im Compute-Kernel hinweg. Eine uniforme Variable kann in einem einzigen skalaren Register gespeichert und von allen Threads im Compute-Kernel wiederverwendet werden.
Affine-Variable
Eine Variable mit Werten, die eine lineare Funktion der Thread-ID für jeden Thread im Compute-Kernel sind. In Algorithmus 11.1 kann die Speicheradresse der Variable y[idx] als affine Transformation der Thread-ID threadIdx.x dargestellt werden:
&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;Diese affine Darstellung kann als Paar von skalaren Werten, einer Basis und einer Schrittweite, gespeichert werden, was deutlich kompakter ist als die vollständig ausgebaute Vektor-Darstellung.
__global__ void vsadd( int y[], int a ) {
i
```Hier ist die deutsche Übersetzung der Markdown-Datei, wobei die Kommentare in den Codeblöcken übersetzt wurden:
int idx = threadIdx.x; y[idx] = y[idx] + a; if ( y[idx] > THRESHOLD ) y[idx] = Y_MAX_VALUE; }
Algorithmus 11.1: Beispiel für skalare und affine Operationen in einem Compute-Kernel (aus [Kim et al., 2013]).
Es gibt mehrere Forschungsvorschläge, wie man uniforme oder affine Variablen in GPUs erkennen und ausnutzen kann. Der Rest dieses Kapitels fasst diese Vorschläge in diesen beiden Aspekten zusammen.
## Erkennung von uniformen oder affinen Variablen
Es gibt zwei Hauptansätze, um das Vorhandensein von uniformen oder affinen Variablen in einem GPU-Compute-Kernel zu erkennen: Compiler-gesteuerte Erkennung und Erkennung über Hardware.
### Compiler-gesteuerte Erkennung
Eine Möglichkeit, das Vorhandensein von uniformen oder affinen Variablen in einem GPU-Compute-Kernel zu erkennen, ist, dies über eine spezielle Compiler-Analyse durchzuführen. Dies ist möglich, da die bestehenden GPU-Programmiermodelle CUDA und OpenCL dem Programmierer bereits Möglichkeiten bieten, eine Variable als konstant über den gesamten Compute-Kernel hinweg zu deklarieren, sowie eine spezielle Variable für die Thread-ID. Der Compiler kann eine Kontrollabhängigkeitsanalyse durchführen, um Variablen zu erkennen, die rein von Konstanten und Thread-IDs abhängig sind, und sie als uniform/affin markieren. Operationen, die ausschließlich auf uniform/affinen Variablen arbeiten, sind dann Kandidaten für Skalierung.
AMD GCN [AMD, 2012] verlässt sich auf den Compiler, um uniforme Variablen und skalare Operationen zu erkennen, die von einem dedizierten Skalar-Prozessor gespeichert und verarbeitet werden können.
Asanovic et al. [2013] führen eine kombinierte konvergente und variantenbasierte Analyse ein, die es dem Compiler ermöglicht, Operationen in einem beliebigen Compute-Kernel zu bestimmen, die für Skalierung und/oder affine Transformation geeignet sind. Anweisungen innerhalb der konvergenten Regionen eines Compute-Kernels können in skalare/affine Anweisungen umgewandelt werden. An jedem Übergang von divergenten zu konvergenten Regionen eines Compute-Kernels fügt der Compiler eine `syncwarp`-Anweisung ein, um durch Kontrollfluss induzierte Registerabhängigkeiten zwischen den beiden Regionen zu behandeln. Asanovic et al. [2013] haben diesen Ansatz übernommen.Hier ist die deutsche Übersetzung der Markdown-Datei. Für den Code wurden nur die Kommentare übersetzt, der Code selbst wurde nicht übersetzt.
Seine Analyse zur Erzeugung von Skalaroperationen für die Temporal-SIMT-Architektur [Keckler et al., 2011, Krashinsky, 2011].
Entkoppelte Affine-Berechnung (DAC) [Wang und Lin, 2017] basiert auf einer ähnlichen Compiler-Analyse, um Skalar- und Affine-Kandidaten zu extrahieren, die in eine separate Warp entkoppelt werden sollen. Wang und Lin [2017] erweitern den Prozess mit einer divergenten Affine-Analyse, mit dem Ziel, Stränge von Anweisungen zu extrahieren