HipKittens: schneller und leistungsstarker AMD-Kernel

• HipKittens ist ein Projekt, das hochperformante Kernel und C++-basierte Programmierprimitiven für AMD-GPUs bereitstellt, um die Effizienz von AI-Berechnungen zu erhöhen
• Bestehende Bestandteile des AMD-Ökosystems wie AITER, PyTorch, Triton, TileLang, Composable Kernel zeigen Grenzen durch instabile Performance und unausgereifte Unterstützung
• HipKittens stellt tile-basierte Abstraktion (tile abstraction) in den Mittelpunkt und trennt hardwarespezifische Implementierungen, während eine gemeinsame Schnittstelle zwischen NVIDIA und AMD erhalten bleibt
• Mit in weniger als rund 500 Zeilen Code geschriebenen Kerneln wird eine höhere Leistung erreicht als mit den bisherigen handgeschriebenen Assembler-Kerneln von AMD
• Das Projekt liefert eine praktische Grundlage, um AI-Berechnungen auf Multi-Silicon-Umgebungen auszuweiten, und zeigt die Möglichkeit eines Übergangs zu einem offenen Hardware-Ökosystem auf

---

Aktueller Stand und Probleme des AMD-GPU-Ökosystems

• AI-Berechnungen wurden bislang vor allem rund um einen einzelnen Hardware-Anbieter weiterentwickelt, doch AMD-GPUs bieten in der jüngsten Generation Rechenleistung und Speicherbandbreite auf Spitzenniveau
  • Beispiel: AMD MI355X OAM mit BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, 288 GB Speicher, 8.0 TB/s Bandbreite
• Wegen des fehlenden ausgereiften Software-Stacks kann diese Leistung in realen AI-Workflows jedoch nicht ausgeschöpft werden
• Zu den wichtigsten Software-Komponenten von AMD gehören AITER, PyTorch, Triton, Mojo, TileLang, Composable Kernel (CK)
  • Die SDPA-Llama-GQA-Backward-Kernel von AITER und PyTorch erreichen jeweils nur 30 % bzw. 24 % der SoTA-Performance
  • Mojo bleibt aufgrund von Bankkonflikten (bank conflict) bei rund 50 % Performance
  • TileLang unterstützt nur bis CDNA3 und erhöht durch fehlende Kernfunktionen und die Abhängigkeit von CK die Komplexität
  • Triton hat Schwierigkeiten bei Register-Lifetime-Tracking und der Optimierung von Speicherzugriffen
• In der Folge müssen die schnellsten AMD-Kernel weiterhin von Experten direkt in Assembler geschrieben werden, was Skalierbarkeit und Wartbarkeit stark einschränkt

Vergleich mit dem CUDA-zentrierten Ökosystem

• In der AI-Community gilt das CUDA-Ökosystem als hohe Eintrittsbarriere (CUDA moat)
• Auch die frühere Entwicklung von NVIDIA-Kerneln erforderte jahrelange Arbeit auf Basis von Low-Level-CUDA/CUTLASS
• In jüngerer Zeit wurde die Entwicklung von NVIDIA-Kerneln durch Fortschritte bei DSLs (domänenspezifischen Sprachen) und AI-gestützten Tools vereinfacht
  • Beispiele: ThunderKittens (TK), CuTe DSL, TinyGrad “tinykittens”, TileLang, Gluon
• Vor diesem Hintergrund wird auch für AMD der Bedarf an neuen Programmierprimitiven untersucht

Designprinzipien von HipKittens

• HipKittens ist die für AMD entwickelte Variante nach ThunderKittens (NVIDIA) und ThunderMittens (Apple Silicon)
• Kernkonzept: tile-basierte Abstraktion (tile abstraction)
  1. Allgemeingültigkeit der tile-Abstraktion – Das auf NVIDIA effektive tile-basierte Rechenmodell lässt sich auch auf AMD natürlich anwenden
  2. Architekturspezifische Backend-Implementierung – Speicherzugriffsmuster und Register-Scheduling werden je nach Hardware unterschiedlich entworfen
  3. Anpassungsfähigkeit der Scheduling-Strategie – Auf AMD CDNA3·CDNA4 ist wave-basiertes Scheduling ineffizient, doch Scheduling auf tile-Ebene bewahrt weiterhin Einfachheit und Leistung
• Durch die Trennung von Schnittstelle (tiles und Operationen) und Implementierung (Hardware-Mapping) wird ein gemeinsam auf verschiedene GPU-Architekturen anwendbares Modell vorgeschlagen

Performance-Ergebnisse

• Attention-Forward-Kernel: in etwa 500 Zeilen Code geschrieben und mit besserer Performance als der Assembler-Kernel von AITER
  • Überlegen bei verschiedenen Head-Dimensionen (D) und Sequenzlängen (N) sowohl in Causal- als auch Non-Causal-Konfigurationen
• GEMM-Kernel: mit einer Kernschleife von unter 100 Zeilen aufgebaut und mit Spitzenleistung sowohl bei BF16- als auch FP8-Berechnungen
• Auch bei Attention Backward, Rotary und Fused Dropout-Residual-LayerNorm-Kerneln wurden Verbesserungen gegenüber der bisherigen Bestleistung bestätigt
• Alle Kernel behalten Lesbarkeit und leichte Änderbarkeit bei und erreichen dennoch Performance auf dem Niveau handgeschriebenen Assemblers

Erweiterung hin zu Multi-Silicon-AI

• Um das Potenzial von AI auszuschöpfen, ist die Nutzung vielfältiger und offener Hardware notwendig
• HipKittens will AMD-GPUs zu einer für AI-Entwickler praktisch zugänglichen Plattform machen
• In Kombination mit dem Open-Source-Software-Stack von AMD zeigt das Projekt die Möglichkeit eines Übergangs zu einem AI-Ökosystem auf Multi-Silicon-Basis auf
• Ein Folgebeitrag behandelt die technischen Detailstrukturen von HipKittens (Ankündigung von Part Two)

[GN⁺]

Hacker-News-Kommentar

• Wir haben einen Vertrag mit AMD und trainieren Llama 405B auf dem MI350X für MLPerf.
  Die Lage bei AMD verbessert sich eindeutig. Wenn man eine AMD-GPU hat, würde ich empfehlen, auf pytorch.org auf Linux+ROCm zu klicken und PyTorch zu installieren. Vor drei Jahren war es noch hoffnungslos, aber inzwischen funktionieren die meisten Mainline-Features gut. Ich habe nanochat auf dem MI300X ausgeführt, und es lief einfach sofort. Auch der MI350X ist entsprechend stabil.
  Natürlich liegt AMD noch hinter NVIDIA zurück, und es braucht viel mehr Investitionen in das Software-Ökosystem sowie Compiler und Treiber. Aber verglichen mit der hoffnungslosen Situation vor zwei Jahren sieht man jetzt Hoffnung.
  Wer sehen will, wo dem LLVM-Backend von AMD noch etwas fehlt, sollte den HipKittens-Code mit CUDA Kittens vergleichen.
  Für Training liegen NVIDIA und Google auf Platz 1, AMD auf Platz 2, und danach kommt fast nichts mehr. Intel und Tenstorrent sind noch weit entfernt, bei Huawei sind die Beispiele mit Segfaults abgestürzt, Groq hat den Chipverkauf aufgegeben, Cerebras ist überhaupt nicht zu bekommen. Bei Trainium hat es fünf Tage gedauert, eine Instanz zu bekommen, also habe ich das Interesse verloren.

  • Ich frage mich, wie weit Tinygrad noch davon entfernt ist, solche Speicheroptimierungen oder Warp-spezifischen Funktionen ausdrücken oder erkunden zu können. Mich würde auch interessieren, wie komplex es wäre, solche Funktionen in tinygrad einzubauen.
  • Ich frage mich, ob man für ROCm + PyTorch auf Consumer-Hardware (nicht MI) ebenfalls einen proprietären Kernel-Treiber braucht.
  • Die Aussage „Cerebras ist nirgendwo zu bekommen“ klingt eher wie ein Zeichen des Sieges.
  • Ich arbeite seit zwei Jahren als CEO einer AMD NeoCloud. Daher freut es mich sehr, diesen Turnaround von AMD direkt mitzuerleben.
    Die Ersteinrichtung ist immer noch etwas holprig, aber verglichen mit früher ist es deutlich besser geworden. Noch vor einem Monat lief nanochat zum Beispiel nicht richtig, jetzt schon. Wichtig ist, dass die Leute sich nun für das AMD-Ökosystem interessieren.
    KI braucht Hardware-Vielfalt. Wenn ein einziges Unternehmen die gesamte KI-Hardware und -Software monopolisiert, mag das gut für Aktionäre sein, aber es schadet dem technischen Fortschritt.

• Ich verstehe die Marktbewertung von NVIDIA nicht. Im Moment haben einige wenige Algorithmen wie Transformer gewonnen, und Daten sind wichtiger geworden. Anders als früher braucht man nicht mehr so viele unterschiedliche HPC-Codes, daher müssen Wettbewerber jetzt nur noch eine engere Menge von Algorithmen optimieren.
  Wenn man vLLM und Transformer effizient ausführen kann, öffnet sich ein riesiger Markt. Dann müssten AMD oder Huawei doch eigentlich leicht aufholen können. Ich frage mich also, was der eigentliche Burggraben (moat) von NVIDIA ist. Reicht InfiniBand allein aus?

  • Ja, der Burggraben von NVIDIA wird immer schwächer. Große Unternehmen wie MS, Google, Amazon und Apple entwickeln alle eigene Chips, um die Abhängigkeit von NVIDIA zu vermeiden.
    Bei Rechenzentrums-GPUs ist NVIDIA zwar weiterhin stark, aber Google nutzt TPUs im großen Maßstab, und auch OpenAI bestellt AMD-Hardware.
    Das CUDA-Ökosystem ist weiterhin wichtig, aber es gibt deutliche Bewegungen, sich davon zu lösen. AMD, Intel, Qualcomm und andere steigen ebenfalls in diesen Wettbewerb ein. HipKittens wirkt wie ein wichtiger Schritt hin zu einem neutralen Software-Ökosystem.
  • Der einfachste Weg, „einige wenige Algorithmen“ zu implementieren, ist der Bau von Allzweck-Rechenhardware. Hardware hat lange Entwicklungszyklen, deshalb ist dieser Ansatz robust.
  • InfiniBand wird durch UEC ersetzt. Für Inference braucht man kein InfiniBand. Deshalb gibt es im Inference-Markt keinen Burggraben. AMD oder Google TPU zu verwenden, ist vernünftig.
  • Die eigentliche Waffe von NVIDIA ist das riesige CUDA-Bibliotheksökosystem. Es gibt Code für nahezu jeden Bereich.
  • Transformer ist nicht nur eine einzige Technologie. Die Implementierungen sind sehr vielfältig. Deshalb sind vLLM oder TRL auch nicht so einfach.

• Man könnte meinen, AMD würde solche Projekte finanzieren, aber ich vermute, sie haben wieder einmal eine Chance verpasst. Bei GPU und KI war es immer so.

  • AMD investiert nur in die minimale Software, die nötig ist, um Produkte auszuliefern. Selbst das Performance-Testsystem ist schwach, und Regressionen werden unverändert an Kunden weitergereicht.
    HipKittens ist eine Verbesserung, aber intern fehlt AMD die Fähigkeit, Kernel-Performance nachzuverfolgen. DevOps wurde an TCS in Indien ausgelagert, und dort kennt man die Lage nicht gut.
    Teams mit guten Führungskräften betreiben auf eigene Faust ein Schatten-IT-Team. ROCm hatte so ein Team nicht, und erst als große Cloud-Kunden protestierten, begannen die Verbesserungen.
    Selbst bei Neueinstellungen werden unter dem Markt liegende Gehälter angeboten. Der Vergleichsmaßstab liegt eher auf dem Niveau von Qualcomm oder Walmart.
    In den letzten vier Jahren wurde der Bonus nie vollständig ausgezahlt.
  • Der Spruch „AMD verpasst keine Gelegenheit, Gelegenheiten zu verpassen“ passt perfekt. Die Instinct-Hardware ist praktisch konkurrenzfähig zu NVIDIA, aber die Software-Unterstützung war furchtbar.
    Wie früher bei Radeon VII wurde der Support eingestellt oder das Ökosystem immer wieder umgeworfen, sodass nichts ausreifen konnte. Wegen mangelnder CUDA-Kompatibilität und fehlender Investitionen haben die meisten Organisationen einfach NVIDIA gewählt.
    Trotzdem investiert AMD in letzter Zeit kontinuierlich in das ROCm- und Instinct-Ökosystem, sodass es sich Stück für Stück verbessert. Im Netzwerkbereich liegt NVIDIA aber weiterhin deutlich vorne.
  • Wenn man sich Leistungsdiagramme ansieht, hätte AMD inzwischen auf NVIDIA-Niveau sein können. Gescheitert ist es aber am Mangel an Software. Chipdesign ist die schwierigere Aufgabe, und trotzdem war das eigentliche Problem, dass man Software nicht verstanden hat.
  • Früher gab es Leute, die für ROCm optimierte Kernel beigetragen haben, aber AMD hat die PRs geschlossen und nicht gemergt. Wirklich unverständlich.
  • Inzwischen ist das Projekt finanziert und funktioniert ordnungsgemäß.

• Ich frage mich, ob es Projekte gibt, die auf ThunderKittens aufbauen.
  Wenn das eine auf HIP portierte Version ist, dann hätte es wohl einen viel größeren praktischen Nutzen als ein einfacher CUDA-Port.

• Die Formulierung „raw assembly vs cooked assembly“ ist interessant.
  Früher war es üblich, Assembler-Code für CPUs direkt zu schreiben. Bei GPUs muss man davor nicht zu viel Angst haben. Am Ende erzeugt der Compiler diesen Code sowieso.

• Mathematisch gesehen ist Inference grundlegende Lineare Algebra bzw. BLAS.
  Ich frage mich, ob mit Blick auf dtype und Sparsity eine kompakte Inference-API möglich wäre, die 80 % der Anwendungsfälle abdeckt. Sie müsste wohl nicht so komplex sein wie CUDA.

• composable-kernel ist die Software, die auf meinem Gentoo-System am häufigsten OOM (Out of Memory) ausgelöst hat. Wenn Clang CK kompiliert, verbraucht es mehr als 2,5 GB pro Thread.

  • Ich habe mir das CK-Paket für Debian angesehen, und beim Build mit -j32 gab es selbst auf einer 64-GB-Workstation OOM. Mit -j1 hat es dann nach 190 Stunden funktioniert.
    Auf dem offiziellen Build-Server sollte man wohl die Zahl der Architekturen reduzieren. Soweit ich gehört habe, brauchen die meisten Abhängigkeitspakete nur Header. Es laufen Arbeiten, um die Build-Zeit zu verkürzen.
  • Dass die Template-Instanziierung für einen einzelnen Kernel mehr als 10 Minuten dauert, ist wirklich erstaunlich.
    device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance — ich habe keine Ahnung, was man Clang da antut.

• Ich frage mich, ob solche Fortschritte dazu führen werden, dass sich auch auf Consumer-AMD-Chips LLMs gut ausführen lassen.
  Ich ziehe zum Beispiel ein HP-OMEN-MAX-16-Notebook mit AMD-CPU/iGPU und RTX 5080 in Betracht. Könnte die AMD-Seite mit der RTX konkurrieren?

  • Man könnte denken, dass eine dGPU immer schneller ist, aber die Begrenzung des Speichervolumens bremst sie aus.
    Dieser Blogbeitrag oder die fortgeschrittenen Mac-Ergebnisse sind interessant.
  • Ich lasse Qwen3 Coder 30B mit Ollama auf einer RTX7900XTX laufen. Es funktioniert ziemlich gut. Ein Teil der Last scheint auf den Systemspeicher und eine Ryzen-7-CPU auszuweichen.
    Es ist etwas langsamer als die Sonnet-4-API, aber für diese Leistung lokal bin ich völlig zufrieden.
    Die Kombination aus Opencode + Ollama + Qwen3 Coder ist eine großartige Alternative zu ClaudeCode + Sonnet4.
    Wenn KI allerdings das gesamte Coding übernehmen soll, sieht man das vielleicht anders. Als persönliche Unterstützung ist es aber perfekt.

• Ich verstehe nicht, warum AMD den B300 komplett ignoriert hat.