Registerdruck und Spills

Register sind der schnellste Speicher eines Threads, aber jeder SM verfügt nur über eine begrenzte Anzahl davon. Wie intensiv ein Kernel sie verwendet, wird als Registerdruck bezeichnet.

Jeder aktive Thread verwendet eine Registerdatei pro SM. Je mehr Register jeder Thread benötigt, desto weniger Threads können gleichzeitig aktiv bleiben.

Eine hohe Registernutzung begrenzt direkt, wie viele Warps auf einen SM passen. Dieser Rückgang der Auslastung kann dazu führen, dass der Hardware zu wenig Arbeit bleibt, um Latenzen zu verbergen.

Wenn ein Thread mehr Register benötigt, als vorhanden sind, lagert der Compiler zusätzliche Werte aus. Dieser Überlauf ist ein Register-Spill in langsameren Speicher.

Ausgelagerte Werte werden in den lokalen Speicher geschrieben, der sich im externen DRAM befindet. Jeder Spill ersetzt einen schnellen Registerzugriff durch einen langsamen Hin- und Rückweg.

Fordern Sie den Compiler auf, die Nutzung zu melden. Wenn Sie --ptxas-options=-v zu nvcc hinzufügen, werden die Register pro Thread sowie alle Spill-Bytes für jeden Kernel ausgegeben.

Der Bericht listet Spill-Speicherungen und -Ladevorgänge auf. Jeder Wert ungleich null bei spill ist ein Warnzeichen dafür, dass Ihr Kernel für langsamen Datenverkehr im lokalen Speicher bezahlt.

Sie können beim Kompilieren eine Obergrenze festlegen. Das Flag --maxrregcount begrenzt die Register pro Thread und tauscht dadurch etwas Geschwindigkeit gegen eine höhere Auslastung.

Ein Hinweis pro Kernel ist oft besser. __launch_bounds__ teilt dem Compiler Ihre Blockgröße mit, sodass er die Register für die gewünschte Auslastung einplanen kann.

Oft können Sie einfach weniger Werte gleichzeitig halten. Das erneute Berechnen eines günstigen Ergebnisses oder das Einschränken des Gültigkeitsbereichs einer Variablen reduziert die Anzahl der aktiven Register.

ILP und das Aufrollen von Schleifen erhöhen den Druck, während Begrenzungen die Auslastung steigern. Die Kunst besteht darin, das Gleichgewicht zu finden, bei dem der Code am schnellsten läuft – und nur der Profiler kann Ihnen sagen, wo es liegt.

Wohin werden Werte verschoben, wenn einem Kernel die Register ausgehen?

Sie haben gelernt, dass hoher Registerdruck die Auslastung verringert und Spills in den langsamen DRAM verursachen kann. Messen Sie die Nutzung, begrenzen Sie die Register und lassen Sie sich vom Profiler leiten. 🎯