PyTorch Helion kann jetzt NVIDIA-Compiler-Steuerung zusammen mit GPU-Kernel-Einstellungen optimieren

Die Ebene unterhalb des Kernels

Laut dem PyTorch-Blogbeitrag wurde Helion im vergangenen Jahr veröffentlicht, damit Entwickler GPU-Kernel in PyTorch-nativem Python schreiben und die Optimierung dem Autotuner überlassen können. Die Autoren geben an, der Autotuner durchsuche Tausende von Kernel-Konfigurationen und liege regelmäßig nur wenige Prozent hinter von Experten optimiertem Code, übertreffe ihn oft sogar.

Der Beitrag argumentiert, dass auch danach noch Spielraum bleibt. Kernel-Autotuner, ob in Helion, Triton oder torch.compile, durchsuchen immer dieselben Dinge: Tile-Formen, Warp-Anzahlen, Pipeline-Stufen und Speicherzugriffsstrategie. Darunter liegt eine zweite Gruppe von Entscheidungen, nämlich wie die Arbeit in Maschinenbefehle übersetzt wird, die ptxas erzeugt. Die Richtlinie für die Registerzuweisung, das feingranulare Instruction Scheduling und Heuristiken der Codegenerierung bleiben bei allen ausgelieferten Kernels auf ihren Standardwerten. Die Autoren sagen, das liege nicht daran, dass die Standardwerte optimal wären, sondern daran, dass bisher nichts danach gesucht habe.

Diese Compiler-Steuerungsebene stellt ptxas selbst ab CUDA 13.3 bereit. Der Beitrag weist darauf hin, dass ihre Erkundung schwierig ist, weil die Steuerungen sich gegenseitig beeinflussen: Eine Änderung an der Registerzuweisung verschiebt die Occupancy, was wiederum entscheiden kann, ob eine Scheduling-Entscheidung hilft oder schadet. Die richtige Einstellung hängt vom jeweiligen Workload ab.

So funktioniert die gemeinsame Suche in Helion

Helion stellt die Compiler-Steuerungen über zwei Einstellungen bereit. Eine ACF ist ein binäres Artefakt, das beim Kompilieren angewendet wird, portabel und versionierbar ist und sich zusammen mit dem Kernel ausliefern lässt.

• advanced_controls_file fixiert eine bewährte ACF innerhalb einer Konfiguration, die der Nutzer bereits gewählt hat.
• autotune_search_acf nimmt eine Liste von ACFs entgegen und lässt den Autotuner von Helion sie als eine weitere Dimension durchsuchen, neben Tile-Größen und Warp-Anzahlen.

Im Suchmodus fügt Helion der Liste automatisch einen leeren String hinzu, der für keine ACF und die Compiler-Standardwerte steht. Der Standard ist damit immer ein Kandidat, und eine ACF wird nur gewählt, wenn sie bei der gemessenen Laufzeit gewinnt. Laut Beitrag besteht der schlimmste Fall darin, dass etwas Autotuning-Budget verbraucht wird und man den Kernel zurückbekommt, den man ohnehin schon hatte.

Die Autoren argumentieren, dass das Tuning in einer einzigen Schleife dem Tuning in zwei Schritten überlegen ist. Blockgröße und Stufenanzahl bestimmen den Registerdruck und den Shared-Memory-Bedarf, und die Richtlinie für die Registerzuweisung entscheidet, ob dieser Druck zu Spills führt. Wird die Kernel-Konfiguration zuerst festgelegt, wird ein Steuerungssatz, der mit einer anderen Konfiguration gut funktionieren würde, nie ausgewertet. Die umgekehrte Reihenfolge hat dasselbe Problem. Deshalb sitzt die Suche im Autotuner und ist kein nachgelagerter Verarbeitungsschritt.

Der Preis ist die Suchzeit, da jede ACF die Zahl der Kompilier- und Messdurchläufe vervielfacht. Der Beitrag empfiehlt, die Suche einmal offline für einen wichtigen Kernel durchzuführen und das Ergebnis anschließend mit advanced_controls_file zu fixieren, so wie heute schon automatisch optimierte Konfigurationen fixiert werden.

Woher die ACFs kommen

NVIDIAs CompileIQ erzeugt die Dateien. Das Tool behandelt die Auswahl der Compiler-Steuerungen als Hyperparameter-Optimierungsproblem und zieht dafür einen kuratierten Suchraum für den Ziel-Compiler und die Toolkit-Version heran. Anschließend läuft eine evolutionäre Suche: Eine Population von Kandidaten-Steuerungssätzen wird vorgeschlagen, jeder Satz kompiliert, die Korrektheit gegen eine Referenz geprüft, die verbleibenden Kandidaten werden zeitlich vermessen, die Gewinner gekreuzt, und der Vorgang wiederholt sich. Das Ergebnis ist eine ACF.

Der Suchalgorithmus enthält nichts GPU-Spezifisches. Der GPU-spezifische Teil ist die Zielfunktion, die der Nutzer schreibt, in der Regel wenige Zeilen rund um triton.testing.do_bench. Laut Beitrag müssen die meisten Nutzer diese Suche nicht selbst ausführen, weil CompileIQ vortrainierte ACF-Kataloge veröffentlicht, die der Autotuner von Helion neben seinen bestehenden Kernel-Konfigurationen auswerten kann.

Eine bewusst harte Vergleichsbasis

Die Autoren betonen, dass der Ausgangspunkt entscheidend ist. Ihre Vergleichsbasis ist das Verzeichnis pretuned_kernels/ im Helion-Repository mit rund 27 Kernels, die Formen aus echtem Modellcode abdecken: FP8- und NVFP4-Quantisierung, RMSNorm-Varianten, RoPE, Grouped GEMM, kausale Depthwise-Faltung, Gated-DeltaNet-Decode sowie fusionierte Bereiche wie eine Qwen3-Decode-Schicht.

Diese Kernels werden nicht im laufenden Betrieb optimiert. Jeder wird mit vorab erzeugten Heuristikdateien pro Architektur ausgeliefert, zum Beispiel je einer für H100 und B200, die Helion zur Laufzeit auswählt. Die Suche fand offline statt, sodass Nutzer optimierte Leistung ohne Suchaufwand erhalten. Jeder Gewinn darüber hinaus sei, so der Beitrag, echtes Verbesserungspotenzial.

Was das für Entwickler bedeutet

Für Entwickler, die bereits auf Helion setzen, wird Tuning auf Compiler-Ebene zu einer Option innerhalb desselben Autotuning-Workflows. Da der Standard in der Kandidatenliste bleibt, ist der Versuch bei einem leistungskritischen Kernel bis auf die Suchzeit mit wenig Risiko verbunden. Teams, die die Suche nicht selbst ausführen wollen, können stattdessen die veröffentlichten ACF-Kataloge von CompileIQ nutzen.