Streaming-Multiprozessoren, Blöcke und Threads (CUDA)

Question

Was ist die Beziehung zwischen einem CUDA-Kern, ein Streaming-Multiprozessor und dem CUDA-Modell von Blöcken und Threads?

Was wird zu dem, was abgebildet und was parallelisiert ist und wie? und was effizienter ist, maximiert die Anzahl der Blöcke oder die Anzahl der Threads?

---

Mein gegenwärtiges Verständnis ist, dass es 8 CUDA-Cores pro Multi-Prozessor. und dass jeder CUDA Kern in der Lage, einen cuda Block zu einem Zeitpunkt auszuführen. und alle Fäden in diesem Block sind seriell in diesem bestimmten Kern ausgeführt.

Ist das richtig?

Answer 1

Die Faden / Blockanordnung ist im Detail in dem CUDA-Programmierhandbuch . Insbesondere Kapitel 4 heißt es:

  

Die CUDA-Architektur ist ein skalierbares Array von multithreaded Streaming Multiprozessoren herum gebaut (SMs). Wenn ein CUDA Programm auf dem Host-CPU einen Kernel grid aufruft, werden die Blöcke des Rasters aufgezählt und Multiprozessoren mit verfügbarer Ausführungskapazität verteilt ist. Die Gewinde eines Gewindeblocks auszuführen gleichzeitig auf einem Multi-Prozessor und mehrere Gewindeblöcke können gleichzeitig auf einem Multi-Prozessor auszuführen. Als Thread blockiert beenden, neue Blöcke werden auf den frei gewordenen Multiprozessoren ins Leben gerufen.

Jedes SM enthält 8 CUDA Kerne und zu einem beliebigen Zeitpunkt sind sie eine einzige Kette von 32 Fäden Ausführung - so es 4 Taktzyklen dauert einen einzelnen Befehl für die ganze Kette zu erteilen. Sie können, dass die Fäden in jedem annehmen WARP im Gleichschritt gegeben auszuführen, aber zu synchronisieren über Warps, müssen Sie Verwendung __syncthreads().

Answer 2

Für die GTX 970 gibt es 13 Streaming-Multiprozessoren (SM) mit 128 CUDA-Cores je. Cuda Cores sind auch Stream-Prozessoren (SP) genannt.

Sie können die Gitter definieren, die Blöcke auf die GPU-Karten.

Sie können die Blöcke definieren, die Fäden zu Stream-Prozessoren (die 128 Cuda Cores pro SM) abzubilden.

ein Kettfaden 32 wird immer durch Fäden gebildet und alle Fäden einer Kette sind sie simultan ausgeführt.

Um die volle mögliche Leistung einer GPU verwenden Sie brauchen viel mehr Threads pro SM als die SM SPs hat. Für jede Compute Capability gibt es eine bestimmte Anzahl von Fäden, die in einem SM zu einem Zeitpunkt befinden können. Alle Blöcke, die Sie definieren in die Warteschlange gestellt und warten auf eine SM die Ressourcen haben (Anzahl der SPs frei), wird es geladen. Die SM beginnt Warps auszuführen. Da nur eine Warp 32 Threads und eine SM hat beispielsweise 128 SPs kann ein SM 4 Warps zu einem bestimmten Zeitpunkt auszuführen. Die Sache ist, wenn die Fäden Speicherzugriff tun wird der Thread blockiert, bis seine Speicheranforderung erfüllt ist. In Zahlen: Eine arithmetische Berechnung auf der SP hat eine Latenzzeit von 18 bis 22 Zyklen, während ein nicht-zwischengespeicherten Zugriff globaler Speicher kann bis zu 300-400 Zyklen dauern. Dies bedeutet, wenn die Fäden eines Kett- für Daten warten nur eine Teilmenge des 128 SPs arbeiten würde. Dafür Scheduler schaltet eine weitere Kette, falls verfügbar auszuführen. Und wenn diese Kette Blöcke führt sie die nächste und so weiter. Dieses Konzept wird Latenz Versteck genannt. Die Anzahl der Kett- und die Blockgröße bestimmen die Belegung (ab, wie viele Verwerfungen der SM können wählen, ausführen). Wenn die Belegung hoch ist, ist es unwahrscheinlich, dass es keine Arbeit für die SPs ist.

Ihre Aussage, dass jeder CUDA Kern einen Block zu einem Zeitpunkt ausgeführt wird, ist falsch. Wenn Sie sich über Streaming-Multiprozessoren sprechen können sie Warps aus allen Thread ausgeführt, die in der SM befinden. Wenn ein Block mit einer Größe von 256 Threads und GPU allowes 2048 Threads resident pro SM hat, würde jede SM hat 8 Blöcke mit Wohnsitz, aus denen die SM Warps können wählen, auszuführen. Alle Fäden der Kettfäden ausgeführt werden parallel ausgeführt.

Sie finden Zahlen für die verschiedenen Compute-Funktionen und GPU-Architekturen hier: https://en.wikipedia.org/wiki/CUDA#Limitations

Sie können eine Belegung Berechnungsblatt von Nvidia herunterladen Occupancy Berechnungsblatt (von Nvidia) .

Answer 3

Der Vertrieb Compute Die Arbeit wird einen Thread-Block (CTA) auf einem SM plant nur, wenn die SM über ausreichende Mittel für den Thread Block (gemeinsam genutzter Speicher, Warps, Register, Barrieren, ...) haben. Thread-Block-Level-Ressourcen wie zum freigegebenen Speicher zugewiesen. Die Zuweisung für alle Threads in dem Thread-Block ausreichend Kett- erzeugt. Der zuordnet Ressourcenmanager Warps mit Round-Robin zu den SM-Unterteilungen. Jede SM Subpartition enthält einen Warp Scheduler, Registerdatei und Ausführungseinheiten. Sobald eine Kette mit einem Subpartition zugeordnet ist, wird es auf dem Subpartition bleiben, bis sie abgeschlossen ist oder von einem Kontextschalter (Pascal Architektur) vorgegriffen. Auf Kontextschalter wiederherstellen wird die Kette auf die gleichen SM gleiche verzugs ID gestellt werden.

Wenn alle Fäden in Kette haben die Kett-Schedulers für alle ausstehenden Befehle abgeschlossen durch die Kette, um eine vollständige ausgegeben und dann werden die Ressourcen-Manager gibt die warp Ebene Ressourcen, die umfassen Kett-id und der Registerdatei.

Wenn alle Verwerfungen in einem Thread Block vollständig dann Block-Level-Ressourcen freigegeben werden und die SM meldet die Compute Arbeits Händler, dass der Block abgeschlossen hat.

Nachdem eine Kette mit einem Subpartition zugeordnet wird und alle Ressourcen werden die Kette zugeordnet ist aktiv Bedeutung angenommen, dass die Kett-Scheduler aktiv den Zustand der Kette verfolgt. Bei jedem Zyklus Scheduler die Kette zu bestimmen, welche sind aktive Verwerfungen ins Stocken geraten und die sind berechtigt, einen Befehl zu erteilen. Die Kett- Scheduler wählt die höchste Priorität qualifizierten Kett- und Ausgaben 1-2 aufeinanderfolgende Befehle aus der Kette. Die Regeln für die Dual-Issue sind spezifisch für jede Architektur. Wenn eine Kette eine Speicherlast ausgibt kann es ausgeführt unabhängige Befehle fort, bis er eine abhängige Instruktion erreicht. Die Kette wird dann, bis die Last abgeschlossen ist ins Stocken geraten berichten. Das gleiche gilt für abhängige mathematische Anweisungen. Die SM-Architektur ist so ausgelegt pro Zyklus durch Umschalten zwischen Kett- sowohl ALU und Speicherlatenzzeit zu verstecken.

Diese Antwort verwendet nicht den Begriff CUDA Kern als Dies führt ein falsche mentales Modell. CUDA Kerne sind pipeline single precision floating point / Integer-Ausführungseinheiten. Die Emissionsrate und die Abhängigkeit Latenz ist zu jeder Architektur spezifisch. Jedes SM Subpartition und SM hat andere Ausführungseinheiten einschließlich der Lade- / Speichereinheiten, double precision Gleitkommaeinheiten, halb precision floating point Einheiten Verzweigungseinheiten, etc.

Um die Leistung zu maximieren die Entwickler zu verstehen, die Trade-off von Blöcken vs. Warps vs. Register / Thread.

Der Begriff Belegung ist das Verhältnis von aktiven Verwerfungen zu maximal Verwerfungen auf einem SM. Kepler - Pascal Architektur (außer GP100) haben 4 Warp Scheduler pro SM. Die minimale Anzahl von Kettfäden pro SM sollte mindestens der Anzahl der Kett- Schedulern gleich sein. Wenn die Architektur eine abhängige Ausführungslatenz von 6 Zyklen (Maxwell und Pascal) hat dann würde mindestens 6 Kett- pro scheduler benötigt die 24 pro SM ist (24/64 = 37,5% Auslastung), um die Latenzzeit abzudecken. Wenn die Fäden Parallelität auf Befehlsebene haben, dann könnte dies reduziert werden. Fast alle Kerne Ausgabe mit variabler Latenz Anweisungen wie Speicherlasten, die 80 bis 1000 Zyklen dauern kann. Dies erfordert mehr aktive Verwerfungen pro Kett-Scheduler zu verbergen Latenz. Für jeden Kernel gibt es einen Trade-off Punkt zwischen der Anzahl der Kett- und anderen Ressourcen wie Shared Memory oder Register so für 100% Auslastung zu optimieren ist nicht wie einige andere Opfer geraten wird wahrscheinlich gemacht werden. Der Profiler CUDA helfen kann Befehlsausgaberate, Auslastung und Stall Gründe zu identifizieren, die Entwickler zu helfen, dieses Gleichgewicht zu bestimmen.

Die Größe eines Gewindeblock kann die Leistung auswirken. Wenn der Kernel große Blöcke und verwendet Synchronisationsbarrieren hat dann Stände Barriere kann ein Gründe kommen Stall sein. Dies kann durch eine Verringerung der gelindert werdenWarps pro Thread-Block.

Answer 4

Es gibt mehr Streaming-Multiprozessor auf einem Gerät.
Eine SM kann mehrere Blöcke enthalten. Jeder Block kann mehrere Threads enthalten.
Ein SM hat mehrere CUDA-Cores (als Entwickler, sollten Sie darüber nicht kümmern, weil sie von Kett- abstrahiert), die am Faden arbeiten. SM arbeitet stets an Kett- Fäden (immer 32). Ein Kettfaden wird arbeiten nur am Faden von demselben Block.
SM und blockieren beide Grenzen auf der Anzahl der Gewinde haben, die Anzahl der Register und den gemeinsamen Speicher.